cuda: decode fmt=9 (bf16) in the dense matmul

[DebugSultan]

cuda: decode fmt=9 (bf16) in the dense matmul

Why

bf16 is already a first-class dtype on the CPU side of this repo: bf16_to_f32
lives in c/st.h:122 and is used by qwen38_core.h, deepseek_v41.c,
inkling.c and st.h's own tensor readers. The CUDA backend does not know it.
A checkpoint whose dense side is bf16 end to end therefore has to keep every
dense projection on the CPU even when a device is present and the expert tier is
already running on it.

This adds the missing decode: fmt=9 in coli_cuda_matmul.

What it does

• row_bytes(9, I) == I * 2 — two raw bytes per weight, no scale array.
• bf16_at(): bf16 -> f32 is the identity on the top 16 bits, so the decoder
  is one shift. No rounding, no lookup table, and no <cuda_bf16.h>
  dependency — this compiles on the same toolchains everything else here does.
  It mirrors st.h's bf16_to_f32 bit for bit, NaN and Inf patterns included.
• A dedicated branch in quant_matmul, alongside the existing 4/6/7/8 branches.

The shared scale-free predicate

fmt=9 carries its magnitude in the weights, so like fmt=0 (f32) and fmt=6
(E8, scales inside the block) it has no scale array to allocate, upload, charge
against the device budget or refresh. Upstream spells that rule out by hand, and
in more than one wording: fmt && fmt != 6 at seven call sites, a longer variant
in the matmul epilogue, and !fmt || fmt==6 — the same rule read backwards
— in coli_cuda_tensor_update. Adding a fourth scale-free format to nine
hand-written lists spelled two different ways is how such lists drift, so this PR
consolidates them into one predicate:

__host__ __device__ static int fmt_scale_free(int fmt) {
    return fmt == 0 || fmt == 6 || fmt == 9;
}

Every converted site keeps the same truth table for formats 0..8, so no existing
container changes behaviour. Three of the conversions deserve a note:

• coli_cuda_tensor_update() is the one that mattered. It ended in
  return !tensor->fmt || tensor->fmt==6 || cudaMemcpy(tensor->scales, scales, ...).
  That is the scale-free rule written in the opposite direction, so a sweep for
  fmt && fmt != 6 never surfaces it, and fmt=9 fell straight through to a
  memcpy from a NULL host pointer — invalid argument on the first in-place
  refresh. The test below caught it, and that is the concrete reason this PR
  consolidates the rule instead of adding a ninth hand-written list.
• absorb_scale() tested !fmt and would have read wscale[row] for a
  scale-free format. It is unreachable in practice — absorb_fmt_ok() gates the
  absorb kernels on coli_cuda_weight_at_supported(), which admits {0,1,2,3,4}
  only — so this one is hardening, not a fix for a live bug.
• the matmul epilogue already excluded fmt != 6 explicitly; routing it through
  the predicate is what keeps fmt=9 out of partial[0] * scales[o], where
  scales is NULL.

What it deliberately does not do

coli_cuda_weight_at_supported() keeps upstream's {0,1,2,3,4}. Callers ask
that predicate "can weight_at decode it", and bf16 rides coli_cuda_matmul's
own branch instead — the same arrangement fmt=6, 7 and 8 already have. The
header comment now records that choice so a later reader does not "fix" it.

Test

c/tests/test_bf16_cuda.cu, wired into make cuda-test. Four independently
falsifiable claims:

1. DECODE — bf16_at() reproduces st.h's bf16_to_f32 over all 65536 bit
   patterns, compared as raw u32 so a NaN cannot hide a mismatch.
2. PARITY — coli_cuda_matmul(fmt=9) against a double CPU reference on a
   shape with an odd I (tail) and S > 1.
3. SCALE-FREE — upload with scales == NULL, coli_cuda_tensor_bytes == I*O*2, and an in-place coli_cuda_tensor_update() with a NULL scale pointer.
   This claim is not decorative: it failed on the first run, and that failure is
   what exposed the !fmt || fmt==6 site above.
4. NEGATIVE CONTROL — a scaled format (fmt=1) with scales == NULL must
   still be refused, so claim 3 cannot pass by the guard being gone.

It includes backend_cuda.cu directly (kernel-level oracle), the same pattern
test_fp8_cuda.cu and test_cuda_fmt_trap_cuda.cu use, and runs early in the
recipe: it is the cheapest kernel test there and has no host-dependent failure
mode.

🤖 Generated with Claude Code

[kreuzzelg]

Reviewed as the author of the dense path this extends (coli_cuda_matmul staging, qt_dense_*, the grouped fmt 1 in #1552). Approve, with two notes.

The consolidation is the right call. fmt && fmt != 6 spelled nine times in two directions is exactly the kind of list that drifts, and coli_cuda_tensor_update ending in !fmt || fmt==6 || cudaMemcpy(...) is the proof: a sweep for the forward form never finds the backward one. fmt_scale_free() keeps the truth table for 0..8 (I checked every converted site against the old expression) and the NULL-memcpy the test caught is a real fix on its own. bf16_at() as a 16-bit shift mirrors st.h's bf16_to_f32 bit for bit, NaN/Inf included, and needs no <cuda_bf16.h> -- good for the toolchains the DLL job builds with.

1. This line meets #1552 (grouped int8 trunk) in the epilogue. #1552 changes the same statement to

(fmt && fmt != 4 && fmt != 6 && fmt != 7 && fmt != 8 && !(fmt == 1 && gs > 0)) ? partial[0] * scales[o] : partial[0]

because fmt 1 uploaded with a group size takes the grouped branch (scales applied per group inside the loop) and must not be multiplied by scales[o] again. Whichever of the two lands second rebases; the merged form is

(!fmt_scale_free(fmt) && fmt != 4 && fmt != 7 && fmt != 8 && !(fmt == 1 && gs > 0)) ? ...

Flagging it here so the gs > 0 term does not get lost in the conflict -- losing it would double-scale every grouped trunk matrix silently (the selftest would catch it, the CI would not).

2. What I did not do. I have not run test_bf16_cuda.cu here: the two cards are serving Ollama at the moment and the house rule is that Colibrì and Ollama never share the GPUs. The CI has only the syntax check for CUDA, so until someone with a free card runs it, the test's evidence is the author's run. I will run it on the 3070/Quadro pair as soon as the cards are free and post the line.

One small thing, not blocking: the PR body's absorb_scale paragraph is cut off mid-sentence ("would have re"), probably a paste.

[kreuzzelg]

Ran test_bf16_cuda.cu on the RTX 3070 (sm_86) now that the cards are free, with one change needed to build it here:

tests/test_bf16_cuda.cu:56:32: error: taking address of temporary array

coli_cuda_init((int[]){0}, 1) is a C99 compound literal; gcc 13.3 as nvcc 12.0's host compiler refuses to take its address in C++17 (clang accepts it as an extension, which is probably why it built on your side). The other CUDA tests use a named array (test_fp8_cuda.cu:152 coli_cuda_init(devs,1)); int dev0 = 0; coli_cuda_init(&dev0, 1); builds everywhere. With that:

[CUDA] device 0: NVIDIA GeForce RTX 3070, 8.2 GB VRAM, sm_86
  decode: 65536/65536 bf16 patterns bit-exact
  parity: 192 outputs vs the double CPU reference
  accounting: 32896 bytes, no scale buffer charged
  refresh: in-place update with NULL scales accepted
  negative control: fmt=1 + NULL scales still refused
bf16 (fmt=9) oracle: 0 failures

So the evidence stands on a second toolchain and card once the literal is a variable. Approval holds; please fold the one-liner in.

[JustVugg]

One thing before this goes in: the build fix from the review is not on the branch yet. tests/test_bf16_cuda.cu still calls coli_cuda_init((int[]){0}, 1), and gcc as nvcc's host compiler refuses the address of a compound literal in C++ ("taking address of temporary array"). A named array does it:

int devs[1] = {0};
if (!coli_cuda_init(devs, 1)) { printf("FAIL cuda init\n"); return 1; }

Push that and I will merge; the kernel itself and the scale-free predicate are fine, and the sweep over all 65536 patterns against bf16_to_f32 is exactly the evidence wanted.