[laudney]

Cloudflare/clef-flash backbone (Qwen3.5-9B, BF16) on the Radeon AI PRO R9700	Before	After
Publisher's BF16 safetensors	refused: no BF16 projection format	loads as is (no MLX or other conversion)
Final-norm rows, mean |error| against FP32, 245–1,024 tokens	—	0.01242 (transformers BF16: 0.01265)
same, 1,025–4,096 tokens	—	0.00954 (transformers BF16: 0.01002)
same, 4,097–10,240 tokens	—	0.00930 (transformers BF16: 0.00964)
same, 10,241–16,384 tokens	—	0.00901 (transformers BF16: 0.00941)
Teacher-forced last-row top-1 vs transformers BF16 (40 records)	—	40/40 equal
Rows with the same prompt cut, repeated	—	byte-identical
Chat on a BF16 Qwen3.5 checkpoint	—	refused by name: a decision backbone only

Part of #1360, closes #1952.

Why

Clef-Flash's joint-schema decisions run on its publisher's BF16 Qwen3.5-9B backbone. Serving the publisher's own weights, not a re-quantized copy, is what makes a benchmark against vLLM or SGLang meaningful and what lets someone adopt the model without asking which conversion they got. This PR is the backbone only: the readout requests and the head follow in the stacked PRs.

What

• feat(quant): a bf16 projection format beside the MLX affine widths.
• feat(qwen3_5): HF BF16 tensor names, the conv layout and the norms of the release checkpoint.
• feat(hip): BF16 projections and embeddings through rocBLAS (hip.Blas.project), with the engine owning a planned 32 MiB rocBLAS workspace so the startup memory plan stays exact.
• fix(hip): look up the rocBLAS workspace calls on the BF16 path only: Blas.open resolves only the product calls; rocblas_set_workspace, rocblas_get_device_memory_size and rocblas_is_managing_device_memory are resolved when the BF16 engine first sets its workspace, and a missing one is refused by name, so GLM-OCR and MLX, which share the loader, never need them.
• test(qwen35) and tools/clef_backbone_rows.py: dump final-norm rows natively and from transformers in BF16 and FP32 over a seeded 40-record corpus (10 per length bucket), then compare.

Receipts

Gate 1 (rows): the native rows are 2–5% nearer FP32 than transformers' own BF16 rows in every bucket, and the teacher-forced top-1 through transformers' BF16 LM head matches on all 40 records. Single rows are not judged: full-attention layers amplify an input error near a tie of scores, and transformers BF16 has a few rows as far from FP32 as native has.

Gate 2 (seams): rocBLAS picks its Tensile kernel by the row count M, so a row's bits can change between prompt cuts of 512, 1,024 and whole; the same cut is byte-identical across repeats. A decision therefore never depends on the cut: a readout always starts at row 0 and never resumes from a kept state, its pieces are pinned to the 1,024-row chunk whatever else runs (#2047, with a control that fails without the pin), and chat on a BF16 checkpoint is refused, so nothing decodes from a kept prefix on this backbone. Byte-identical seams would need fixed-M calls or a pinned solution (upstream's blas_fixed.zig for GLM-OCR is that route), which is a follow-up if wanted.

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

[ashhart]

Thanks @laudney, this is a clean base for the Clef-Flash stack. The bf16 format stays opt-in by the config alone, and the MLX path is untouched on the host. Refusing to serve until the readout lands is the right call too.
Four small things before it lands:

• a run of an existing MLX Qwen3.5 checkpoint on HIP, byte-equal to main, with the contract SUMMARY for the new embed records;
• the rocBLAS version that has the three new workspace calls, or a lookup on the BF16 path only, since GLM-OCR shares that loader;
• the compare report the receipt table came from;
• a line on how the stack will keep a resumed prompt equal to a fresh one, given that rocBLAS changes bits with the row count.

[laudney]

Ask	Answer
MLX Qwen3.5 on HIP, byte-equal to main	Qwen3.5-9B-4bit (mlx-community 8b2b98c0) on the R9700, /v1/chat/completions, greedy, 128 tokens, 4 prompts: token ids and reply text byte-equal to upstream main (table below)
Contract SUMMARY for the embed records	SUMMARY gfx1201 77 kernels: 77 pass gates 1-4, 0 fail, 2 known findings; 7 controls: 7 red as planted, 0 missed; 2 limit problems. The 2 limit problems are tf_sum_partials and tf_sum_partials_act (recorded 23/24 VGPRs, HIP 7.16 gives 21/22); main has the same two
rocBLAS workspace calls	b347a40 fix(hip): look up the rocBLAS workspace calls on the BF16 path only: Blas.open resolves only the product calls; the three workspace calls are resolved on first use, which only the BF16 engine makes, and a missing one is refused by name. GLM-OCR and MLX never resolve them. The oldest rocBLAS I can check (5.1.1) already has all three
The compare report	below (gate 1, 40 records)
A resumed prompt equal to a fresh one	a readout always starts at row 0 and never resumes from a kept state; its pieces are pinned to the 1,024-row chunk (5b9fa01 in #2047, with a control that fails without the pin), so every rocBLAS call it makes has the same row count whatever else runs; and chat on a BF16 checkpoint is refused, so nothing decodes from a kept prefix on this backbone

Prompt (MLX Qwen3.5-9B-4bit, greedy, 128 tokens)	main	this branch	equal
printing press	5ca97c21fe36c724	5ca97c21fe36c724	yes
Fibonacci code	0929d30afc9fd611	0929d30afc9fd611	yes
garden tips	0787d0e9259ce57d	0787d0e9259ce57d	yes
hash collisions	49b5e04282182860	49b5e04282182860	yes

Embed record lines: tf_embed_dense (bf16 table) and (fp32 table) pass nonfinite, repeat and shapes (the fp32 table bit-equal to the kernel it replaces); both planted controls (gathers the next row; reads a bf16 table as fp16) red at shapes; live limits 256 threads, 13 VGPRs, 0 LDS, as recorded.

gate 1 compare report (tools/clef_backbone_rows.py compare)

{
 "sources": {
  "torch": "2.15.0a0+rocm10.1.0a20260822",
  "transformers": "5.22.1",
  "checkpoint": "Cloudflare/clef-flash@fb4417d31c1d0f9e7e77f1141b40aaede7491e00"
 },
 "buckets": [
  {
   "bucket": "245-1024",
   "records": [
    "rec_00",
    "rec_01",
    "rec_02",
    "rec_03",
    "rec_04",
    "rec_05",
    "rec_06",
    "rec_07",
    "rec_08",
    "rec_09"
   ],
   "rows": 6357,
   "native_mean": 0.012417430112239893,
   "bf16_mean": 0.012646992549515845,
   "passes": true
  },
  {
   "bucket": "1025-4096",
   "records": [
    "rec_10",
    "rec_11",
    "rec_12",
    "rec_13",
    "rec_14",
    "rec_15",
    "rec_16",
    "rec_17",
    "rec_18",
    "rec_19"
   ],
   "rows": 27234,
   "native_mean": 0.009537120340056611,
   "bf16_mean": 0.010023162834097656,
   "passes": true
  },
  {
   "bucket": "4097-10240",
   "records": [
    "rec_20",
    "rec_21",
    "rec_22",
    "rec_23",
    "rec_24",
    "rec_25",
    "rec_26",
    "rec_27",
    "rec_28",
    "rec_29"
   ],
   "rows": 68786,
   "native_mean": 0.009299759542405747,
   "bf16_mean": 0.009638283130318349,
   "passes": true
  },
  {
   "bucket": "10241-16384",
   "records": [
    "rec_30",
    "rec_31",
    "rec_32",
    "rec_33",
    "rec_34",
    "rec_35",
    "rec_36",
    "rec_37",
    "rec_38",
    "rec_39"
   ],
   "rows": 137161,
   "native_mean": 0.00901123333391803,
   "bf16_mean": 0.009409329436653966,
   "passes": true
  }
 ],
 "top1": []
}

[ashhart]

Thanks @laudney, that answers all four.

• The workspace lookup on the BF16 path is exactly what we wanted. GLM-OCR and MLX never touch those calls now, and a missing one is named.
• The greedy run, the contract SUMMARY and the compare report all line up with the table.
• The readout rule (row 0, no resume, 1,024-row pieces, chat refused) is the right answer for resumed against fresh.

Two small things, neither blocking:

• one seeded sampled run of the same four prompts, beside the greedy ones;
• please move the resume rule into the body's gate 2 paragraph, and list the workspace fix under "What", so the body matches the code.

We don't have an AMD GPU, so your posted runs are the device evidence for this.

[laudney]

Prompt (MLX Qwen3.5-9B-4bit, seeded sampling, 128 tokens)	main	this branch	equal
printing press	ad3ab466f8052ca4	ad3ab466f8052ca4	yes
Fibonacci code	78d1380d0b40ade9	78d1380d0b40ade9	yes
garden tips	d1ae49e7954cd664	d1ae49e7954cd664	yes
hash collisions	d056e22ac0290226	d056e22ac0290226	yes

The seeded sampled run beside the greedy one: R9700, main f23603801 against this branch b347a40fd, /v1/chat/completions with "temperature": 0.7, "top_p": 0.9, "seed": 1234, "max_tokens": 128, the same server flags as the greedy run (--parallel 1 --context 4096 --no-thinking --prompt-cache-gib 0, default drafts). Each build ran twice and gave the same bytes both times, and every sampled reply differs from the greedy one of the same prompt, so sampling took effect.

The body now puts the resume rule in the gate 2 paragraph and lists the workspace fix under "What".

[ashhart]

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.