Clef-Flash: joint-schema decisions on /v1/systemone (HIP) by laudney · Pull Request #2048 · ashhart/TensorFold

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

Cloudflare/clef-flash on the Radeon AI PRO R9700	Before	After
`/v1/systemone` with a Clef-Flash checkpoint	refused	the release's answers from its joint schema head on the GPU
End to end vs the reference BF16 pipeline, mean / max KL, 245–1,024 tokens	—	6.9e-05 / 7.9e-04
same, 1,025–4,096 tokens	—	8.9e-05 / 1.5e-03
same, 4,097–10,240 tokens	—	3.7e-05 / 2.8e-04
same, 10,241–16,384 tokens	—	3.4e-05 / 1.8e-04
Top choice agrees with the reference (130 questions, 40 records)	—	130/130
One 16,384-token record, end to end	—	3.9 s
Two 16K readouts and a short request at once: steps between a request's turns	—	≤ 8 (bound 11)
Head finish and release, 16K readout	—	16–19 ms

Stacked on #2047 (readout requests) and #2046 (BF16 backbone). Closes #1360. The encoder is gilby's port from #1361, with his authorship kept and Ash's review asks resolved.

Why

Clef's decisions are only as good as the head that reads the backbone's rows, so the head runs on the GPU from the start, as vLLM and SGLang serve it, over the publisher's own BF16 weights. Serving the publisher's checkpoint unchanged is what makes the numbers above comparable and what lets someone adopt Clef-Flash without a conversion step.

What

feat(server): encode Clef records (gilby) and fix(server): match the Clef encoder to the release: the release's rendering, option order and spans, pinned against the real tokenizer's ids.
feat(joint_schema): the head config, loader and an FP64 host reference.
feat(hip): strided batched rocBLAS products, ten joint schema kernels with kernel contracts and planted controls, and the device head (state per readout stream, scratch planned and counted in the startup budget; refused by name under --tp).
feat(tokenizer) and fix(server): apply pretokenize_regex only to a Clef checkpoint: for a checkpoint that ships joint_head_config.json, tokenizer_config.json's pretokenize_regex replaces tokenizer.json's Split, as transformers 5's Qwen3.5 tokenizer does (Clef-Flash's tokenizer.json still has the older \p{L}+ split; 12 of 12 pinned records match only with the field). Every other checkpoint keeps tokenizer.json's split; applying the field more widely (it also changes LightOnOCR-3) is a separate PR with token-equality receipts.
feat(server): answer System One from the joint schema head: 422 validation errors with their loc, 400 refusals at the design's bounds, model echoed, raw probabilities.
fix(native): refuse a quantized checkpoint that ships the joint head: refused by name before any load.
fix(lanes): pin a readout's pieces to the prompt chunk (from lanes: readout requests for decision prompts #2047) and fix(hip): look up the rocBLAS workspace calls on the BF16 path only (from qwen3_5: serve Clef-Flash's BF16 Qwen3.5 backbone on HIP #2046).
docs(api): the route, its fields, bounds, statuses and answer shape.
feat(lanes): trace engine steps to a file: opt-in TF_LANE_TRACE=<file>, one JSON line per engine-thread action; one null check when unset.
tools/clef_head_gate.py, tools/clef_e2e.py, tools/clef_serving.py: gates 3, 4 and 6.

Receipts

Gate 3 (head alone on the reference's BF16 rows, 367 logits): native equals the reference BF16 head in 341–342 logits; of the rest it is nearer FP64 in 11–12 and farther in 13–15 (no direction). Its worst error ties the reference's; its mean is 0.2–6% higher per type, so the strict "no larger" rule fails by rounding noise. Disclosed, not tuned.
Gate 4 (end to end): the table above, against the rule mean KL ≤ 1e-4 and max KL ≤ 1e-2 per bucket, with no top-choice disagreement. Every answer has the reference's keys, choices and token counts, but none is byte-equal: 477 of 498 rounded values differ, by at most 0.024.
Gate 5 (operators): HIP kernel contract on gfx1201 at this branch: 89 kernels: 89 pass gates 1-4, 0 fail, 2 known findings; 10 controls: 10 red as planted, 0 missed; 2 limit problems; the 2 limit problems are main's tf_sum_partials and tf_sum_partials_act records (23/24 VGPRs recorded, 21/22 under HIP 7.16), not this PR's. The new kernels' 40 launch-limit records (10 kernels × gfx1030, gfx1100, gfx1151, gfx1201) are current for ROCm 10.1's core/bin hipcc (HIP 7.16), the release compiler.
Gate 6 (serving): every readout step is one piece of ≤ 1,024 rows; the completion sits in the last step; a client that leaves frees its lane; the short request's bytes are the same alone and beside the readouts.

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

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

Thanks @laudney for this. The stack splits cleanly: the 13 lower commits match #2047 exactly, so this review covers the 15 on top. Thanks also to @gilby, whose encoder port is credited on its commit.

The head is well separated. /v1/systemone changes only for a BF16 checkpoint that ships joint_head_config.json. Readouts always start at row 0, so a resumed request can't differ from a fresh one.

A few things before it can land:

The pretokenize_regex change reaches every model on every backend, and it changes LightOnOCR-3's tokens. Could it be scoped to Clef, made opt-in, or split into its own PR with token-equality receipts?
lane_host.zig is now 604 lines, so the lean checker is one problem over baseline.
A quantized checkpoint that carries the head file opens as a text model; please refuse it by name.
Please pin readout pieces to 1,024 rows whatever other lanes are doing, with a host test.
Please regenerate the launch records with the release compiler, attach the gate tools' output files, fix the gate 4 wording, and document the route in docs/api.md.

laudney · 2026-11-03T18:06:05Z

Ask	Answer
`pretokenize_regex` reaches every model	`78e04dc` `fix(server): apply pretokenize_regex only to a Clef checkpoint`: the field applies only when the checkpoint ships `joint_head_config.json`; every other checkpoint, LightOnOCR-3 included, keeps tokenizer.json's split (a test covers both sides). Applying it more widely becomes its own PR with token-equality receipts
`lane_host.zig` over 600 lines	`9aab994` `fix(lanes): keep the lane trace portable`: 599 lines; lean against the merge base: 424 problems, 0 new. It also fixes the Windows CI break (the trace used `std.c.open`/`write`; it now uses `std.Io` files), checked by a cross-compile to x86_64-windows-gnu
A quantized checkpoint with the head file	`489c088` `fix(native): refuse a quantized checkpoint that ships the joint head`: "the joint schema head needs the BF16 checkpoint: … is quantized (…)", before any load; the same MLX checkpoint without the head file still opens as a text model
Pin readout pieces to 1,024 rows	`54e6658`, the same commit as #2047's `5b9fa01` (host test, plus a device control that fails without the pin)
Launch records with the release compiler	the 40 joint records (10 kernels × gfx1030/1100/1151/1201) are already current for ROCm 10.1's `core/bin` hipcc (HIP 7.16); no change. gfx1201 contract: `89 kernels: 89 pass gates 1-4, 0 fail, 2 known findings; 10 controls: 10 red as planted, 0 missed; 2 limit problems`; the 2 are main's `tf_sum_partials` records, not this PR's
Gate tools' output files	below
Gate 4 wording	fixed in the PR text
`docs/api.md`	`6a4bc53` `docs(api): document /v1/systemone for a Clef engine`

Also on this branch: 8ed05a7, #2046's rocBLAS workspace lookup. zig build test -Dhip=true on gfx1201 (ROCm 10.1): 527/527 steps.

gate 3: head alone (tools/clef_head_gate.py compare; pieces 512, 1,024, whole)

// piece 512
{
 "sources": {
  "torch": "2.15.0a0+rocm10.1.0a20260822",
  "transformers": "5.22.1",
  "checkpoint": "Cloudflare/clef-flash@fb4417d31c1d0f9e7e77f1141b40aaede7491e00"
 },
 "types": [
  {
   "type": "noul",
   "logits": 117,
   "native_mean": 0.004529234780813532,
   "native_worst": 0.01836959644573355,
   "bf16_mean": 0.004571493435311176,
   "bf16_worst": 0.019342133123873673,
   "pass": true
  },
  {
   "type": "choice",
   "logits": 137,
   "native_mean": 0.005714737139596581,
   "native_worst": 0.019521818010266756,
   "bf16_mean": 0.005619153935056486,
   "bf16_worst": 0.019521818010266756,
   "pass": false
  },
  {
   "type": "score",
   "logits": 113,
   "native_mean": 0.004692881484317049,
   "native_worst": 0.016606180742905208,
   "bf16_mean": 0.004577627515388878,
   "bf16_worst": 0.016606180742905208,
   "pass": false
  }
 ],
 "paired": {
  "logits": 367,
  "identical": 342,
  "native_nearer": 12,
  "native_farther": 13
 },
 "pass": false
}
// piece 1024
{
 "sources": {
  "torch": "2.15.0a0+rocm10.1.0a20260822",
  "transformers": "5.22.1",
  "checkpoint": "Cloudflare/clef-flash@fb4417d31c1d0f9e7e77f1141b40aaede7491e00"
 },
 "types": [
  {
   "type": "noul",
   "logits": 117,
   "native_mean": 0.004635338047557729,
   "native_worst": 0.019342133123873673,
   "bf16_mean": 0.004571493435311176,
   "bf16_worst": 0.019342133123873673,
   "pass": false
  },
  {
   "type": "choice",
   "logits": 137,
   "native_mean": 0.0058654373197163595,
   "native_worst": 0.019521818010266756,
   "bf16_mean": 0.005619153935056486,
   "bf16_worst": 0.019521818010266756,
   "pass": false
  },
  {
   "type": "score",
   "logits": 113,
   "native_mean": 0.004589175820600234,
   "native_worst": 0.016606180742905208,
   "bf16_mean": 0.004577627515388878,
   "bf16_worst": 0.016606180742905208,
   "pass": false
  }
 ],
 "paired": {
  "logits": 367,
  "identical": 341,
  "native_nearer": 12,
  "native_farther": 14
 },
 "pass": false
}
// piece whole
{
 "sources": {
  "torch": "2.15.0a0+rocm10.1.0a20260822",
  "transformers": "5.22.1",
  "checkpoint": "Cloudflare/clef-flash@fb4417d31c1d0f9e7e77f1141b40aaede7491e00"
 },
 "types": [
  {
   "type": "noul",
   "logits": 117,
   "native_mean": 0.004609585124538492,
   "native_worst": 0.019342133123873673,
   "bf16_mean": 0.004571493435311176,
   "bf16_worst": 0.019342133123873673,
   "pass": false
  },
  {
   "type": "choice",
   "logits": 137,
   "native_mean": 0.0058519057049638,
   "native_worst": 0.019521818010266756,
   "bf16_mean": 0.005619153935056486,
   "bf16_worst": 0.019521818010266756,
   "pass": false
  },
  {
   "type": "score",
   "logits": 113,
   "native_mean": 0.0048575382457218385,
   "native_worst": 0.016606180742905208,
   "bf16_mean": 0.004577627515388878,
   "bf16_worst": 0.016606180742905208,
   "pass": false
  }
 ],
 "paired": {
  "logits": 367,
  "identical": 341,
  "native_nearer": 11,
  "native_farther": 15
 },
 "pass": false
}

gate 4: end to end (tools/clef_e2e.py compare)

{
 "sources": {
  "torch": "2.15.0a0+rocm10.1.0a20260822",
  "transformers": "5.22.1",
  "checkpoint": "Cloudflare/clef-flash@fb4417d31c1d0f9e7e77f1141b40aaede7491e00"
 },
 "buckets": [
  {
   "bucket": "245-1024",
   "records": 10,
   "questions": 36,
   "mean_kl": 6.894260182840258e-05,
   "max_kl": 0.0007868636980129899,
   "top_agree": 36,
   "pass": true
  },
  {
   "bucket": "1025-4096",
   "records": 10,
   "questions": 29,
   "mean_kl": 8.893350446908835e-05,
   "max_kl": 0.0014877856793438237,
   "top_agree": 29,
   "pass": true
  },
  {
   "bucket": "4097-10240",
   "records": 10,
   "questions": 29,
   "mean_kl": 3.691297999426149e-05,
   "max_kl": 0.0002838000869669049,
   "top_agree": 29,
   "pass": true
  },
  {
   "bucket": "10241-16384",
   "records": 10,
   "questions": 36,
   "mean_kl": 3.4323175294531606e-05,
   "max_kl": 0.0001752600500723768,
   "top_agree": 36,
   "pass": true
  }
 ],
 "disagreements": [],
 "answers_equal": 0,
 "records": 40,
 "lengths": [
  {
   "record": "rec_00",
   "target": 245,
   "actual": 245
  },
  {
   "record": "rec_01",
   "target": 256,
   "actual": 256
  },
  {
   "record": "rec_02",
   "target": 544,
   "actual": 582
  },
  {
   "record": "rec_03",
   "target": 600,
   "actual": 600
  },
  {
   "record": "rec_04",
   "target": 612,
   "actual": 612
  },
  {
   "record": "rec_05",
   "target": 650,
   "actual": 650
  },
  {
   "record": "rec_06",
   "target": 661,
   "actual": 661
  },
  {
   "record": "rec_07",
   "target": 785,
   "actual": 785
  },
  {
   "record": "rec_08",
   "target": 980,
   "actual": 980
  },
  {
   "record": "rec_09",
   "target": 1024,
   "actual": 1024
  },
  {
   "record": "rec_10",
   "target": 1025,
   "actual": 1025
  },
  {
   "record": "rec_11",
   "target": 1109,
   "actual": 1109
  },
  {
   "record": "rec_12",
   "target": 1551,
   "actual": 1551
  },
  {
   "record": "rec_13",
   "target": 1764,
   "actual": 1764
  },
  {
   "record": "rec_14",
   "target": 2447,
   "actual": 2447
  },
  {
   "record": "rec_15",
   "target": 3385,
   "actual": 3385
  },
  {
   "record": "rec_16",
   "target": 3780,
   "actual": 3780
  },
  {
   "record": "rec_17",
   "target": 4021,
   "actual": 4021
  },
  {
   "record": "rec_18",
   "target": 4056,
   "actual": 4056
  },
  {
   "record": "rec_19",
   "target": 4096,
   "actual": 4096
  },
  {
   "record": "rec_20",
   "target": 4097,
   "actual": 4097
  },
  {
   "record": "rec_21",
   "target": 4707,
   "actual": 4707
  },
  {
   "record": "rec_22",
   "target": 4846,
   "actual": 4846
  },
  {
   "record": "rec_23",
   "target": 5156,
   "actual": 5156
  },
  {
   "record": "rec_24",
   "target": 6500,
   "actual": 6500
  },
  {
   "record": "rec_25",
   "target": 7024,
   "actual": 7024
  },
  {
   "record": "rec_26",
   "target": 8173,
   "actual": 8173
  },
  {
   "record": "rec_27",
   "target": 8832,
   "actual": 8832
  },
  {
   "record": "rec_28",
   "target": 9211,
   "actual": 9211
  },
  {
   "record": "rec_29",
   "target": 10240,
   "actual": 10240
  },
  {
   "record": "rec_30",
   "target": 10241,
   "actual": 10241
  },
  {
   "record": "rec_31",
   "target": 12272,
   "actual": 12272
  },
  {
   "record": "rec_32",
   "target": 12363,
   "actual": 12363
  },
  {
   "record": "rec_33",
   "target": 13040,
   "actual": 13040
  },
  {
   "record": "rec_34",
   "target": 13200,
   "actual": 13200
  },
  {
   "record": "rec_35",
   "target": 13424,
   "actual": 13424
  },
  {
   "record": "rec_36",
   "target": 14679,
   "actual": 14679
  },
  {
   "record": "rec_37",
   "target": 15528,
   "actual": 15528
  },
  {
   "record": "rec_38",
   "target": 16030,
   "actual": 16030
  },
  {
   "record": "rec_39",
   "target": 16384,
   "actual": 16384
  }
 ],
 "pass": true
}

gate 6: serving (tools/clef_serving.py)

{
 "requests": [
  {
   "id": 2,
   "rows": 16384,
   "rules": {
    "lines": true,
    "pieces": 16,
    "ok_one_piece_each": true,
    "contiguous": true,
    "completion_in_last_step": true,
    "submit_before_first_call": true,
    "most_steps_between_turns": 8,
    "bound": 11,
    "within_bound": true
   },
   "t_piece_ms": {
    "mean": 231.5498484,
    "max": 274.472488
   },
   "t_last_piece_and_finish_ms": 297.275766,
   "completion_ms": 17.051393,
   "longest_wait_ms": 2018.516809,
   "submit_to_done_ms": 7509.063838,
   "name": "rec_39"
  },
  {
   "id": 3,
   "rows": 16030,
   "rules": {
    "lines": true,
    "pieces": 16,
    "ok_one_piece_each": true,
    "contiguous": true,
    "completion_in_last_step": true,
    "submit_before_first_call": true,
    "most_steps_between_turns": 2,
    "bound": 11,
    "within_bound": true
   },
   "t_piece_ms": {
    "mean": 230.93334713333334,
    "max": 272.542527
   },
   "t_last_piece_and_finish_ms": 212.941655,
   "completion_ms": 14.780917,
   "longest_wait_ms": 217.82872,
   "submit_to_done_ms": 4289.633322,
   "name": "rec_38"
  },
  {
   "id": 4,
   "rows": 245,
   "rules": {
    "lines": true,
    "pieces": 1,
    "ok_one_piece_each": true,
    "contiguous": true,
    "completion_in_last_step": true,
    "submit_before_first_call": true,
    "most_steps_between_turns": 1,
    "bound": 11,
    "within_bound": true
   },
   "t_piece_ms": {
    "mean": 0.0,
    "max": 0.0
   },
   "t_last_piece_and_finish_ms": 61.132439,
   "completion_ms": 2.114825,
   "longest_wait_ms": 112.43163,
   "submit_to_done_ms": 173.531068,
   "name": "rec_00"
  }
 ],
 "beside_equals_alone": true,
 "leave_cancelled_mid_prompt": true,
 "client_closed": true,
 "after_served": true,
 "stray_failed_lines": [],
 "lanes_in_flight_after": true,
 "pass": true
}

gate 5: kernel contract on gfx1201

# PR #2048: joint schema launch records against the release compiler
Compiler: ROCm 10.1.0 hipcc (HIP 7.16.26385) from ROCm's core/bin. Card for the gates: AMD Radeon AI PRO R9700 (gfx1201).

## zig build hip-kernel-contract-limits -Dhip=true -Dhipcc=<ROCm 10.1>/core/bin/hipcc (gfx1030, gfx1100, gfx1151, gfx1201; no GPU)
joint schema launches: 40 of 40 PASS (record current), 0 stale; no joint record changes
PASS limits gfx1030 joint schema layer norm tf_js_layer_norm: 20 VGPRs, 22 SGPRs, 36 LDS bytes, 0 scratch bytes, bound 256, wave32 (record current)
PASS limits gfx1100 joint schema layer norm tf_js_layer_norm: 20 VGPRs, 22 SGPRs, 36 LDS bytes, 0 scratch bytes, bound 256, wave32 (record current)
PASS limits gfx1151 joint schema layer norm tf_js_layer_norm: 19 VGPRs, 22 SGPRs, 36 LDS bytes, 0 scratch bytes, bound 256, wave32 (record current)
PASS limits gfx1201 joint schema layer norm tf_js_layer_norm: 19 VGPRs, 22 SGPRs, 36 LDS bytes, 0 scratch bytes, bound 256, wave32 (record current)
PASS limits gfx1030 joint schema GELU tf_js_gelu: 8 VGPRs, 10 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record current)
PASS limits gfx1100 joint schema GELU tf_js_gelu: 8 VGPRs, 18 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record current)
PASS limits gfx1151 joint schema GELU tf_js_gelu: 8 VGPRs, 10 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record current)
PASS limits gfx1201 joint schema GELU tf_js_gelu: 8 VGPRs, 7 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record current)
PASS limits gfx1030 joint schema row softmax tf_js_softmax_rows: 19 VGPRs, 15 SGPRs, 72 LDS bytes, 0 scratch bytes, bound 256, wave32 (record current)
PASS limits gfx1100 joint schema row softmax tf_js_softmax_rows: 19 VGPRs, 18 SGPRs, 72 LDS bytes, 0 scratch bytes, bound 256, wave32 (record current)
PASS limits gfx1151 joint schema row softmax tf_js_softmax_rows: 19 VGPRs, 15 SGPRs, 72 LDS bytes, 0 scratch bytes, bound 256, wave32 (record current)
PASS limits gfx1201 joint schema row softmax tf_js_softmax_rows: 19 VGPRs, 14 SGPRs, 72 LDS bytes, 0 scratch bytes, bound 256, wave32 (record current)
PASS limits gfx1030 joint schema widening tf_js_widen: 4 VGPRs, 9 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record current)
PASS limits gfx1100 joint schema widening tf_js_widen: 5 VGPRs, 18 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record current)
PASS limits gfx1151 joint schema widening tf_js_widen: 4 VGPRs, 8 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record current)
PASS limits gfx1201 joint schema widening tf_js_widen: 4 VGPRs, 7 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record current)
PASS limits gfx1030 joint schema rounding tf_js_round: 5 VGPRs, 9 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record current)
PASS limits gfx1100 joint schema rounding tf_js_round: 4 VGPRs, 18 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record current)
PASS limits gfx1151 joint schema rounding tf_js_round: 4 VGPRs, 8 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record current)
PASS limits gfx1201 joint schema rounding tf_js_round: 4 VGPRs, 7 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record current)
PASS limits gfx1030 joint schema span sums tf_js_span_sums: 7 VGPRs, 22 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record current)
PASS limits gfx1100 joint schema span sums tf_js_span_sums: 7 VGPRs, 22 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record current)
PASS limits gfx1151 joint schema span sums tf_js_span_sums: 7 VGPRs, 22 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record current)
PASS limits gfx1201 joint schema span sums tf_js_span_sums: 7 VGPRs, 22 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record current)
PASS limits gfx1030 joint schema span means tf_js_span_means: 9 VGPRs, 14 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record current)
PASS limits gfx1100 joint schema span means tf_js_span_means: 9 VGPRs, 18 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record current)
PASS limits gfx1151 joint schema span means tf_js_span_means: 8 VGPRs, 14 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record current)
PASS limits gfx1201 joint schema span means tf_js_span_means: 8 VGPRs, 10 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record current)
PASS limits gfx1030 joint schema option summary tf_js_option_summary: 23 VGPRs, 34 SGPRs, 72 LDS bytes, 0 scratch bytes, bound 256, wave32 (record current)
PASS limits gfx1100 joint schema option summary tf_js_option_summary: 23 VGPRs, 34 SGPRs, 72 LDS bytes, 0 scratch bytes, bound 256, wave32 (record current)
PASS limits gfx1151 joint schema option summary tf_js_option_summary: 23 VGPRs, 34 SGPRs, 72 LDS bytes, 0 scratch bytes, bound 256, wave32 (record current)
PASS limits gfx1201 joint schema option summary tf_js_option_summary: 23 VGPRs, 34 SGPRs, 72 LDS bytes, 0 scratch bytes, bound 256, wave32 (record current)
PASS limits gfx1030 joint schema logits tf_js_logits: 39 VGPRs, 33 SGPRs, 36 LDS bytes, 0 scratch bytes, bound 256, wave32 (record current)
PASS limits gfx1100 joint schema logits tf_js_logits: 38 VGPRs, 33 SGPRs, 36 LDS bytes, 0 scratch bytes, bound 256, wave32 (record current)
PASS limits gfx1151 joint schema logits tf_js_logits: 37 VGPRs, 34 SGPRs, 36 LDS bytes, 0 scratch bytes, bound 256, wave32 (record current)
PASS limits gfx1201 joint schema logits tf_js_logits: 37 VGPRs, 41 SGPRs, 36 LDS bytes, 0 scratch bytes, bound 256, wave32 (record current)
PASS limits gfx1030 joint schema gated logits tf_js_combine: 8 VGPRs, 14 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record current)
PASS limits gfx1100 joint schema gated logits tf_js_combine: 7 VGPRs, 18 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record current)
PASS limits gfx1151 joint schema gated logits tf_js_combine: 7 VGPRs, 14 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record current)
PASS limits gfx1201 joint schema gated logits tf_js_combine: 7 VGPRs, 14 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record current)
Stale records outside this PR (kernels it does not touch; the step exits 1 on them):
FAIL limits gfx1030 planned walk tf_sum_partials: 21 VGPRs, 20 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record stale)
FAIL limits gfx1100 planned walk tf_sum_partials: 21 VGPRs, 20 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record stale)
FAIL limits gfx1151 planned walk tf_sum_partials: 21 VGPRs, 20 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record stale)
FAIL limits gfx1201 planned walk tf_sum_partials: 21 VGPRs, 22 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record stale)
FAIL limits gfx1030 moe shared gate tf_moe_combine_gate: 9 VGPRs, 22 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record stale)
FAIL limits gfx1100 moe shared gate tf_moe_combine_gate: 10 VGPRs, 22 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record stale)
FAIL limits gfx1151 moe shared gate tf_moe_combine_gate: 9 VGPRs, 22 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record stale)
FAIL limits gfx1030 planned walk tf_sum_partials_act: 22 VGPRs, 22 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record stale)
FAIL limits gfx1100 planned walk tf_sum_partials_act: 22 VGPRs, 22 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record stale)
FAIL limits gfx1151 planned walk tf_sum_partials_act: 22 VGPRs, 22 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record stale)
FAIL limits gfx1201 planned walk tf_sum_partials_act: 22 VGPRs, 24 SGPRs, 0 LDS bytes, 0 scratch bytes, bound 1024, wave32 (record stale)

## zig build hip-kernel-contract -Dhip=true -Dhip-arch=gfx1201 -Dhipcc=<ROCm 10.1>/core/bin/hipcc (on the GPU)
SUMMARY gfx1201 89 kernels: 89 pass gates 1-4, 0 fail, 2 known findings; 10 controls: 10 red as planted, 0 missed; 2 limit problems; gates 90207 ms, limits 1404 ms
joint schema kernels: 52 PASS lines, 0 FAIL lines (6 N/A: no sums declared)
The 2 limit problems (outside this PR):
FAIL limits live planned walk tf_sum_partials: 256 threads launched, the runtime allows 1024; 21 VGPRs, 0 + 0 LDS bytes (recorded 23, 0)
FAIL limits live planned walk tf_sum_partials_act: 256 threads launched, the runtime allows 1024; 22 VGPRs, 0 + 0 LDS bytes (recorded 24, 0)

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

Merged #2047's new tip in 820ee37, which brings main and the same hostInfo fix; the CHANGELOG keeps the Clef-Flash entry above main's new ones. Lean against main: 0 new; zig build test -Dhip=true on the R9700 (ROCm 10.1): 532/532 steps.

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

Thanks @laudney, every point is answered:

the regex now applies only with joint_head_config.json, tested both ways;
lane_host.zig is back under 600, with the Windows trace fix on top;
a quantized checkpoint with the head is refused by name;
the readout pin, the docs and the gate outputs are all here.

Gate 3's strict rule failing by rounding noise is disclosed clearly. We'll make that call on our side.

One thing for landing: the export carries #2047's pin commit twice and its rocBLAS commit in two versions, so the series doesn't apply as sent. Once #2047 lands, a rebase onto it will drop them. This lands after #2046 and #2047.

laudney · 2026-11-04T09:55:23Z

Rebased onto #2047's rebased tip (33936299c): 19 commits, with #2047's pin commit and the rocBLAS commit gone, since both are below it now. Main moved under it in four places: joint_schema imports bf16_round as main's module; blas.zig is main's file plus the strided batched product only, byte-equal to the earlier tip's; contract_controls.hip, docs/api.md (Rerank, then System One) and the CHANGELOG keep both sides. Checks on the R9700 (ROCm 10.1): zig build test -Dhip=true 552/552 steps; lean 0 new; HIP kernel contract SUMMARY gfx1201 96 kernels: 96 pass gates 1-4, 0 fail, 2 known findings; 11 controls: 11 red as planted, 0 missed; 2 limit problems (main's tf_sum_partials pair), with no joint schema launch record stale.

ashhart · 2026-11-04T14:36:03Z

Thanks @laudney. Everything from our last review is in, and the rebase onto #2047 dropped the duplicated commits. This lands right after #2047. By then main will have moved in three registration lists (hip_records.zig, hip/root.zig and the HIP contract's main), so one more rebase then, with the SUMMARY line again. We'll give our answer on gate 3's rounding-noise result before it lands.

laudney · 2026-11-04T18:29:55Z

Rebased onto main (9d6f1d78c) after #2047: 20 commits. The three registration lists (hip_records.zig, hip/root.zig, contract_main.zig) keep main's entries first and add the joint schema ones; the CHANGELOG keeps the Clef-Flash entry first. Main had grown lane_host.zig and hip_records.zig to 600 lines, so a new top commit, refactor(lanes): keep lane_host and hip_records within 600 lines, moves three trace comments onto their statements, the trace test import into lane_host_tests.zig and the joint records import onto the line before it; nothing runs differently.

Checks on the R9700 (ROCm 10.1, core/bin hipcc): zig build test -Dhip=true 571/571 steps (7309/7372 tests, 63 skipped); lean against main 425 problems, 0 new; test-lane-api-host 494/494 after the line fix.

SUMMARY gfx1201 107 kernels: 107 pass gates 1-4, 0 fail, 2 known findings; 15 controls: 15 red as planted, 0 missed; 2 limit problems; gates 1452882 ms, limits 1387 ms

The 2 limit problems are main's tf_sum_partials and tf_sum_partials_act; no joint schema launch record is stale.

ashhart · 2026-11-04T20:49:45Z

Thanks @laudney, the rebase after #2047 is clean, and the SUMMARY on the new head is all green apart from main's own two limit problems. On gate 3: we accept the rounding-noise result as you disclosed it. Gate 4's KL and the 130 of 130 top choices are what this route needs, and every answer is exact against itself. This is ready to land.

ashhart · 2026-11-05T09:02:21Z

@laudney, a rebase is owed here, with one thing for it. This conflicts with main, and core/lane_host.zig is at exactly 600 lines on main, so the trace hooks take it to 611 and lean fails. The hooks themselves are fine: off by default, one null check per round, fill call and submit, no scheduling change. The clean fix is to move fillAt whole into lane_host_fill.zig, where fill logic belongs. That frees about 27 lines, with a one-line forwarder if callers stay. It goes back in the queue once CI is green on that head.

laudney · 2026-11-05T09:58:20Z

Rebased onto main (9200a3064): 21 commits.

Ask / main change	Answer
`lane_host.zig` over 600	`refactor(lanes): move fillAt into lane_host_fill.zig`: `fillAt` moves whole, as a file-level function beside its only caller (the fill step), so no forwarder; `firstCommitted` becomes public for its completion hook. `lane_host.zig` 584 lines, `lane_host_fill.zig` 563
`hip_records.zig` at 600	`style(contract)`: the joint records' import shares the line before it, as the other record files' do
#2139 dropped `bf16` from the Qwen3.5 HIP offers while BF16 was refused	with the head a BF16 checkpoint serves, so `bf16` is offered again (both kinds, as `decisionOnly` admits both); the test now checks it is offered, `test-qwen-hip-formats-host` filters the renamed test, and the CHANGELOG entry says so
`build/hip.zig` groups are now `Source`s	the joint group is `.{ .path = "joint_schema/joint_schema.hip" }`, default flags like the others
the records count test (182 / 170)	192 / 180 with the joint schema's 10, in the commit that adds them

Checks on the R9700 (ROCm 10.1, core/bin hipcc): lean 425 problems, 0 new; zig build test -Dhip=true green but the records count, then test-cuda-host 313/313 with the count fixed; lane host steps 573/573.

SUMMARY gfx1201 107 kernels: 107 pass gates 1-4, 0 fail, 2 known findings; 15 controls: 15 red as planted, 0 missed; 2 limit problems; gates 1415688 ms, limits 1374 ms

The 2 limit problems are main's tf_sum_partials pair.

laudney · 2026-11-05T10:39:44Z

Two commits on top (no rebase): test(hip-contract): register the joint schema head as a family and a one-line comment fix. The joint schema kernels now form the joint_schema contract entry (contract_joint_schema.zig), as GLM-OCR does, so contract_main.zig no longer carries their block and hip_records.zig is byte-equal to main's again; their 10 records move into the family list (pinned counts: 192 in all, 170 shared). Each case is seeded by its kernel's name, so the cases are the same as before the move.

Run on the R9700 (ROCm 10.1, `core/bin` hipcc)	Result
`tf-hip-kernel-contract image joint_schema`	`SUMMARY joint_schema HIP full: 12 kernels: 12 pass gates 1-4, 0 fail, 0 known findings; 3 controls: 3 red as planted, 0 missed; 0 limit problems; gates 1273 ms, limits 9 ms`, exit 0
full run	`SUMMARY gfx1201 107 kernels: 107 pass gates 1-4, 0 fail, 2 known findings; 15 controls: 15 red as planted, 0 missed; 2 limit problems` (main's `tf_sum_partials` pair), the same totals as before
host	`check-hip-kernel-contract-host` and `test-cuda-host` 318/318 steps; lean 0 new

laudney · 2026-11-05T11:24:14Z

Rebased onto main (0311c9280, 19 commits on since this morning): 24 commits. tokenizer_test.zig keeps main's Qwen2 slow-tokenizer tests and adds the pretokenize_regex ones after them. Main grew hf_text.zig to 555 lines, so the Clef override's 45-line test took it to 612; a new top commit, refactor(server): move the pretokenize_regex test into its own file, moves it whole into hf_text_regex_test.zig (registered beside hf_eos_test.zig) and hf_text.zig is 568 lines. test-server-cpu runs the same 1,076 tests, and a planted failure in the moved test turns it red.

Checks on the R9700 (ROCm 10.1, core/bin hipcc): lean 425 problems, 0 new; zig build test -Dhip=true 588/588 steps (7497/7560 tests, 63 skipped);

SUMMARY joint_schema HIP full: 12 kernels: 12 pass gates 1-4, 0 fail, 0 known findings; 3 controls: 3 red as planted, 0 missed; 0 limit problems
SUMMARY gfx1201 107 kernels: 107 pass gates 1-4, 0 fail, 2 known findings; 15 controls: 15 red as planted, 0 missed; 2 limit problems

image joint_schema exits 0; the full run's 2 limit problems are main's tf_sum_partials pair.

ashhart · 2026-11-05T12:19:16Z

Thanks @laudney, for keeping hip_records.zig byte-equal to main by moving the joint records into the family list, and for the full SUMMARY on the new head. One thing: capabilities now lists bf16 for HIP Qwen 3.5 again, but its chat server still refuses a BF16 checkpoint, which is what #2139 fixed. Could the offer say it's for decisions only (with joint_head_config.json), so a client doesn't read it as BF16 chat? The lane-host hooks have had their read and are fine as behaviour: off by default, with no scheduling change. Moving fillAt keeps core/lane_host.zig under its limit.

laudney · 2026-11-05T12:43:55Z

One commit on top (no rebase): fix(hip): list BF16 Qwen3.5 under decisions, not chat formats.

`capabilities --json`, HIP	main	before (`075783dab`)	now (`602c72a9f`)
`families.qwen3_5` / `qwen3_5_moe`	the 18 `mlx-q{bits}g{group}`	the 18 + `bf16`	the 18 (main's list, byte-equal `mlx_formats`)
`decisions`	—	—	`{"qwen3_5": ["bf16"], "qwen3_5_moe": ["bf16"]}`

A family now declares decision_formats apart from formats: the formats it serves only for decisions (a readout head beside the checkpoint), never for chat. api.Family carries them (default empty, so CUDA, Metal and XPU write an empty decisions), hip_families.list copies them to each model_type and alias, and cli_capabilities.zig writes them under the new key. A client that reads families as chat formats keeps #2139's answer, and one that wants System One reads decisions. docs/api.md names the key; the CHANGELOG line says the same.

Tests: the capability audit adds a case (decision formats appear under decisions only; an engine with none writes {}), the alias test carries them, and the Qwen3.5 offer test checks bf16 is absent from formats and present in decision_formats for both model types.

Checks on the R9700 (ROCm 10.1): zig build test -Dhip=true 588/588 steps, 7,498 passed, 63 skipped; zig fmt --check; lean on this head merged into main 4a20028aa: 0 new problems. The merge still conflicts only in CHANGELOG.md.

ashhart · 2026-11-05T13:12:20Z

Thanks @laudney. A separate decisions key, with the chat formats byte-equal to main's, keeps #2139's answer and still tells a System One client what it can load. Ready to land after the 1.2.3 cut.

The encoder now takes validated fields whose criteria already have their type's shape (noul overrides, a non-empty choice object or score list), tokenizes each piece through the server's model_text.Text, and sorts object keys in the request arena rather than page_allocator. Hermetic tests on a tiny greedy tokenizer cover the pieces and spans, both option orders, the six instruction fallbacks, null descriptions, numbering from 1, truncation at the limit and Unicode states, and tools/clef_tokens.py pins the reference's ids and spans for twelve records, which a test compares when TF_CLEF_DIR is set.

Turn the joint schema head's BF16 logits into the reference's systemone_answer body: each question's logits widened to FP32 and softmaxed in FP32, answers keyed in request order with the first best choice winning a tie, score sums in double over the levels, and Python's round(x, 4) done exactly on the double's binary value in a u128 scaled integer with ties to even. Non-finite logits fail the body, and "raw" returns the unrounded probabilities in encoding order.

Add core/joint_schema, the host side of the Clef joint schema head: the config with every field checked, the tensor walk that names all 122 tensors with exact shapes, a loader that refuses a missing, extra, non-BF16 or misshaped tensor by name through core/safetensors, and a reference of JointSchemaHead.forward in FP64 or rounding to BF16 at the head's points. tools/clef_head_golden.py runs a generated tiny head through the release's own code on the CPU; the FP64 reference matches its FP64 logits to 1e-14 and the near-zero normalize and cosine cases bit for bit, and the release's head loads when TF_CLEF_DIR is set.

The joint schema head attends over up to 16,384 memory rows with 16 heads, so blas.zig gains rocblas_gemm_strided_batched_ex and two products over the heads: scores, FP32 Q_h K_h^T times the scale with the keys of a query row contiguous for the row softmax, and context, FP32 P_h V_h written at each head's columns. extent takes an element size and FP32 buffers must sit on 4 bytes. The device checks compare both products with float64 at keys 245 to 16,384 and query blocks of 1 to 64, keep their guards, show FP32 outputs, and require a wrong head stride and a transposed layout to fail the same bound.

The joint schema head's BF16 rounding points between rocBLAS products need their own kernels, in a twelfth source group resolved by hip/joint_schema/launch.zig: affine layer norm, exact GELU, an FP32 row softmax, BF16 widening and rounding, span sums in prompt row order over piece rows or table rows by id and their means, step 5's option summary, and step 7's prior, cosine and scorer rows and the gated logit. Each kernel passes contract gates 1 to 4 against float64 at the head's points (rows.zig gains layer norm and GELU), two planted controls go red, docs name them, and gate 5 records each launch.

The qwen3_5 HIP engine now serves a Clef checkpoint: with joint_head_config.json beside a BF16 backbone it loads the checked head, uploads its weights with the model's, plans each readout stream's state (memory rows, span sums, last row) and one shared scratch, and gives the lane backend a Readout. hip/joint_schema runs each piece through LN_hidden, the memory projection and the span sums, and the finish through steps 2 to 7 with 64-row FP32 attention blocks; Readout.open takes the prompt for the lexical rows and copies its ids. A BF16 checkpoint without the head is still refused as decision-only.

tf-clef-test tiny writes tools/clef_head_golden.py's generated head as a checkpoint, loads it through the host loader and the device head, and consumes the 48 rows whole and in pieces of 16, 7 and 1. The logits of every cut must be bit-equal, and within 4 BF16 steps of both torch's BF16 head and the host's BF16 emulation; it reports each contender's worst and mean error against the fixture's FP64 logits. The build compiles it in `zig build test` against stub kernels.

Gate 3 judges the device head alone on the reference's BF16 hidden rows. tools/clef_head_gate.py rebuilds each gate record's spans from clef_backbone_rows.py's seeded schema text by tokenizer offsets, runs the release's JointSchemaHead in BF16 on the GPU and in FP64 on the CPU over a table of the record's own lm_head rows, drives the native head through tf-clef-test head at a given piece size, and compares each contender's mean and worst error against FP64 per question type.

transformers 5's Qwen3.5 tokenizer splits words with its own PRETOKENIZE_REGEX, which keeps combining marks in their words, and tokenizer_config.json publishes that regex as pretokenize_regex; Clef-Flash and LightOnOCR-3 still ship the older \p{L}+ split in tokenizer.json, while Qwen3.5-9B, Qwen3.6-35B-A3B and Qwen3.8-27B already agree and do not change. HfText now replaces the first Split pattern with the field whenever it is present and refuses a field that is empty, not a string or does not compile, naming it. Hermetic tests split a decomposed word both ways, and the pinned Clef goldens match 12 of 12 records through the override.

A readout engine (a Clef checkpoint's joint schema head) now answers /v1/systemone: a missing or mistyped field is a 422 with FastAPI's detail list and its location, an unknown type, empty criteria, more than 64 questions, 256 options a question or 1,024 in all, a score outside 2 to 10 levels, images, videos and a schema past the window of min(16,384, context) are 400s, and a valid record is encoded, submitted untaught as a readout and waited on through the shared prompt-only wait. Its BF16 logits become the reference's answers, `model` echoed, and a non-finite logit or failed run is a 500.

The encoder, the answers and the host and device heads port Cloudflare's Apache-2.0 joint_schema_model.py, so THIRD_PARTY_NOTICES gains the Cloudflare Clef section from ashhart#1361 with the files' new places, the tools that run that file and what their fixtures hold, and the changelog's Unreleased list notes Clef-Flash on /v1/systemone and the pretokenize_regex split, crediting gilby's first encoder port.

Gate 4 judges a served engine's /v1/systemone against the reference's BF16 pipeline. tools/clef_e2e.py writes 40 seeded System One requests cut to the gate's bucket lengths through encode_record, runs load_release_model on the GPU for each record's probabilities and answers, sends the requests with "raw": true to the server, and reports per bucket the mean and worst KL and the top-choice agreement, listing each disagreement with both contenders' top two and gaps against the proposed thresholds.

Gate 6 judges a readout's engine-thread steps, which nothing logged. TF_LANE_TRACE=<file>, read once when the lane host starts, appends one JSON line per action on the awake clock: each submit (written before the host's lock), each prompt call with its rows, each decode round and each readout completion, every line in one write to an append-only file so no thread shares a lock with the engine thread. Unset, each path costs one null check. HIP's prompt step stores a readout's fill_at for the rows; a readout has no drafts and HIP reports no fill, so no reader changes. A host test traces two readouts of three pieces each.

Gate 6 judges how readouts share the engine thread. With the server writing TF_LANE_TRACE, tools/clef_serving.py sends two 16K readouts and a short one at once and checks from the trace that each prompt call is one piece of at most 1,024 rows, contiguous, the last holding the completion, that no calls overlap, and that at most eight plus the lane count other steps run between a readout's turns; it reports T_piece, the last piece with the finish, the completion and each request's longest wait, then checks that a client leaving inside a 16K readout frees its lane by the short request's unchanged bytes.

tokenizer_config.json's pretokenize_regex reached every model that loads through HfText, so it changed how other checkpoints that ship the field, such as LightOnOCR-3, split their words. The loader now applies it only when the checkpoint directory holds joint_head_config.json, the file that makes it a Clef checkpoint; other checkpoints keep tokenizer.json's Split, and the field is ignored there even when it is malformed. The hermetic test covers both sides, and the 12 pinned Clef records still match.

The HIP engine loads the joint schema head only over the plain BF16 weights, so a quantized Qwen3.5 checkpoint that carries joint_head_config.json opened as a text model and quietly dropped the head. The open now refuses it by name ("the joint schema head needs the BF16 checkpoint", with its quantization) before any weight loads, through the same check that refuses a bare BF16 backbone; the host test covers the quantized checkpoint with and without the head.

The lane trace opened and wrote its file with std.c.open and write, which do not compile on Windows, where a file is a HANDLE, so the Windows CI build failed. The trace now creates its file with std.Io.Dir.createFile and writes each line with one streaming write, where the OS keeps a shared offset whole between threads. Its open, close and fill-line helpers move into lane_trace.zig, so lane_host.zig is back under the 600-line limit.

A Clef checkpoint answers /v1/systemone from its joint schema head, and the route's request, answers and refusals were documented nowhere. The new section in docs/api.md gives the request fields (state, questions, their types and criteria, raw), the answer shape for each type and the raw extension, the bounds on questions, options and record length, the FastAPI 422 body and the 400, 503 and 500 refusals, in the style of the Decisions section beside it.

Main brought `lane_host.zig` to its 600-line limit, so the lane trace's hooks took it over; `fillAt` is fill logic and its only caller is the fill step, so it moves whole into `lane_host_fill.zig` as a file-level function, with `firstCommitted` made public for its completion hook, and the host file drops to 584 lines with nothing running differently.

Main brought `hip_records.zig` to its 600-line limit, so the joint schema records' import, which had wrapped onto a line of its own, now shares the line before it, as the other record files' imports do; the records are unchanged.

The joint schema kernels ran inline in the full contract run, so their block in contract_main.zig conflicted with every contract main added, and the run's exit status mixed their result with main's own limit problems. They now form the joint_schema entry, which opens its own runtime, controls and launcher and runs the same 12 kernels and 3 controls with the same name-seeded cases, and their 10 gate-5 records move from the shared list into the family list, so `image joint_schema` checks the head alone and the pinned counts read 192 in all and 170 shared.

The new entry's file comment ran past the 120-column limit that the lean check holds every comment to, so it now names the kernels and their planted controls in one shorter line; nothing else changes.

Main grew `hf_text.zig` toward its 600-line limit, so the Clef override's 45-line test took it over; the test moves whole into `hf_text_regex_test.zig`, registered beside `hf_eos_test.zig` in the file's test block, and `hf_text.zig` drops to 568 lines. The server step runs the same 1,076 tests, and a planted failure in the moved test turns it red, so the test still runs.

The HIP Qwen3.5 servers refuse a BF16 checkpoint for chat and serve it only for /v1/systemone with a joint schema head, so the capabilities list must not show BF16 beside the formats that chat. A family now declares decision_formats apart from formats; api.Family carries them, the HIP list copies them to each model_type and alias, and capabilities --json writes them under a new decisions key. The Qwen3.5 text list is main's again, and docs/api.md names the key.

ashhart · 2026-11-06T02:59:52Z

Thank you @laudney. This is merged into main with your commits under your own name, so you're listed as a TensorFold contributor. It ships in the next release, credited to you in the changelog and the release notes.

This was referenced Nov 3, 2026

Cloudflare Clef decisions on native CUDA: /v1/systemone through Clef's joint schema head (Clef-Flash 9B passes) #1360

Closed

Qwen3.5 BF16 on HIP: the publisher's checkpoint as a decision backbone (Clef-Flash), prompt rows through rocBLAS #1952

Closed

BlivionIaG mentioned this pull request Nov 3, 2026

lfm2: DSpark drafting with owned context and confidence scheduling #1394

Open

5 of 8 tasks

laudney force-pushed the pr-clef-3-joint-head branch from 820ee37 to b879dfd Compare November 4, 2026 09:55

laudney mentioned this pull request Nov 4, 2026

feat(hip): one host protocol for the HIP OCR families #2040

Merged

jschmied mentioned this pull request Nov 4, 2026

kv ledger: count the paged pool's promises in the lane host's ledger, no behaviour change (#1943 PR 1) #2127

Merged

laudney force-pushed the pr-clef-3-joint-head branch from b879dfd to cdb7a52 Compare November 4, 2026 18:29

BobClawblaw mentioned this pull request Nov 4, 2026

Unlimited-OCR-2 on CUDA: the family's steps on NVIDIA GPUs, the router's long-context speedup (stacked on #2074) #2077

Open

8 tasks done

ElderOrb mentioned this pull request Nov 4, 2026

hip: build and load the runtime on Windows #2178

Open

4 tasks done

AbrahamCete mentioned this pull request Nov 5, 2026

test(hip): the planned walk's tf_sum_partials records hold ROCm 10.1's hipcc (HIP 7.16) #2233

Merged

3 of 4 tasks

laudney force-pushed the pr-clef-3-joint-head branch from cdb7a52 to a85018f Compare November 5, 2026 09:58

Bizuayeu mentioned this pull request Nov 5, 2026

glm cuda: TF_KV_SHARE keeps GLM's prompts in the shared prompt store, and GLM's own store goes #2211

Merged

10 tasks done

laudney force-pushed the pr-clef-3-joint-head branch from 09787a6 to 075783d Compare November 5, 2026 11:24

AbrahamCete mentioned this pull request Nov 5, 2026

perf(hip): the decode tail issues a row's loads before its branches #2247

Open

8 tasks done

laudney mentioned this pull request Nov 5, 2026

Nemotron Parse on HIP: NVIDIA's checkpoint, token-exact against its HF example #1941

Open

laudney and others added 25 commits November 6, 2026 02:59

style(contract): keep hip_records.zig within 600 lines

1a2d879

Main brought `hip_records.zig` to its 600-line limit, so the joint schema records' import, which had wrapped onto a line of its own, now shares the line before it, as the other record files' imports do; the records are unchanged.

style(hip-contract): keep the joint schema entry's header in 120

d22eef2

The new entry's file comment ran past the 120-column limit that the lean check holds every comment to, so it now names the kernels and their planted controls in one shorter line; nothing else changes.

changelog: back under Unreleased

fc96840

ashhart force-pushed the pr-clef-3-joint-head branch from 3e64170 to fc96840 Compare November 6, 2026 02:59

ashhart merged commit a757420 into ashhart:main Nov 6, 2026
3 of 7 checks passed

ashhart mentioned this pull request Nov 6, 2026

Changelog: tonight's lines cite their own PRs and thank their authors #2312

Merged

laudney mentioned this pull request Nov 6, 2026

Clef-Flash on native CUDA: /v1/systemone through the joint schema head (GB10) #2324

Open

Conversation

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

Uh oh!

Why

What

Receipts

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!

laudney commented Nov 4, 2026

Uh oh!

ashhart commented Nov 4, 2026

Uh oh!

laudney commented Nov 4, 2026

Uh oh!

ashhart commented Nov 4, 2026

Uh oh!

ashhart commented Nov 5, 2026

Uh oh!

laudney commented Nov 5, 2026

Uh oh!

laudney commented Nov 5, 2026

Uh oh!

laudney commented Nov 5, 2026

Uh oh!

ashhart commented Nov 5, 2026

Uh oh!

laudney commented Nov 5, 2026

Uh oh!

ashhart commented Nov 5, 2026

Uh oh!

Uh oh!

ashhart commented Nov 6, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

laudney commented Nov 3, 2026 •

edited

Loading