cuda: tree attention and DFlash block attention for every family, checked against float64

[jschmied]

What this changes

CUDA has no shared attention for drafted rows. A family that drafts a tree on CUDA writes its own attention over the
committed keys plus each row's path, and a DFlash-style drafter writes its own block attention. This adds both as
shared modules, operator-only (no model calls them yet), next to cuda/tree_round (#734), which does the verify and
commit step of the same round:

• zig/kernels/cuda/tree_attention.cu, zig/src/cuda/tree_attention.zig: drafted rows of any number of streams in
  one call. Each row sees its stream's committed keys and then its root-to-row path (paths() builds the paths from
  tree-local parents, depth up to 128). Keys run in 512-key groups on tensor cores, merged in group order, at head
  dims 128 and 256 with grouped KV heads. A stream's fully committed groups are computed once for all its rows (16
  (row, head) pairs as one mma tile, the same tiles and merge as a row's own block), so every partial keeps its bits
  whether a row runs alone or beside others. The grid is fixed: at most 16 group slots per (row, key head), each
  walking its groups, so a captured graph does not grow with the context. A group's key tiles go to 4 slices; a
  four-warp block combines its slices in order; where four-warp blocks would leave SMs idle (TreeAttention.sms),
  each slice takes its own one-warp block and the merge combines the 4 in the same order (a small fixed slice area
  in the scratch), so the bits do not depend on the device's size. A one-warp block fetches its tiles 2 ahead (3
  cp.async stages where the one-warp grid fits one wave of the SMs, else 1 ahead in 2, two blocks an SM; stages
  never change bits) while it computes one. Partials hold the query heads only. The merge runs a block per (row, key
  head, head).
• zig/kernels/cuda/block_attention.cu, zig/src/cuda/block_attention.zig: a DFlash-style drafter's block rows
  over their stream's context (each stream's K and V cache through a pointer table, within a sliding window) and then the block itself,
  causal or not. One warp per head on tensor cores.
• zig/src/core/attention_ref.zig: one query head's attention in float64 with a per-element tolerance (2^-7 of
  sum p|v| / sum p, plus a bf16 half-ulp), the reference both tests check against.

Users we know of: the Qwen3.8-27B family in #608 (target tree attention at dim 256 and its DFlash2 drafter), and any
CUDA family that drafts trees with #734.

Receipt

• Environment: TensorFold c36b536d (1.0.9) plus these commits; one DGX Spark (GB10, sm_121, CUDA 13). No
  checkpoint for the tests.
• Exactness:
  • tf-cuda-test tree: dim 128 (48 query / 4 KV heads) and 256 (24 / 4), several streams with different committed
    lengths and trees, every element within the float64 tolerance (worst 0.47 of it). A stream's rows run alone
    equal their bits in the batch; one shared block walking every (row, group), a grid sized for a 262k context,
    and one-warp blocks equal the default grid.
  • tf-cuda-test block: windowed contexts then the block, causal and not, block widths 7 and 16, within
    tolerance (worst 0.38). A stream run alone equals its bits in the batch.
• Mutations: each of these fails its test: paths cut at depth 3, the score scale off by 1/16, the window bound
  off by one, the causal mask off by one, the merge leaving out a group's fourth slice (19x the tolerance, and the chunked merge's fourth slice), a tile fetched into the stage being read.
• In a model (fork only, not in this PR): with Qwen3.8-27B on CUDA: the qwen27 family's engine for NVIDIA GPUs, every stream drafted #608's 27B engine wired to these kernels instead of its Triton
  attention, drafted long prompt: attention 46.1 ms against Triton's 49.0 ms; 16 streams at 32k context, code and
  chat at temperature 0 and 1: every bench_concurrent --alone --serial cell exact, throughput within -4..+2 % of
  Triton. Replies equal 0.6.6's except two at bf16 near-ties (an exact top-2 tie, and one ulp), where drafted still
  equals serial. Teacher-forced on 1,023 rows of real text: NLL 0.8058 vs 0.8029, argmax agreement 99.12 %.
  The slices, one-warp blocks and stages (dc54a0e5, 1adfc970, b7d1bdc0) change no bits: 4,092 teacher-forced
  logit rows byte-identical to the build before, drafted and 8-stream shared rounds with the same shas and round
  counts. 512-key groups (01969f6d) give the 27B new bits once, on every device: drafted == serial on all four
  prompts in both block forms, and against the Triton set on 4,092 teacher-forced rows dNLL +0.0009, argmax
  agreement 99.32 %, mean KL 0.00066 (2,048-key groups: +0.0011, 99.12 %, 0.00071). Serial decode after a 4,722-token
  prompt, 128 tokens, GB10: attention 275.8 ms with the first version, 209.1 ms now (one-warp blocks; four-warp
  217.5 ms; Triton 244.3 ms). On an RTX 5090 (Qwen3.8-27B on CUDA: the qwen27 family's engine for NVIDIA GPUs, every stream drafted #608): attention 34.6 µs a call against Triton's 59.7, long 2.92 vs
  2.94 s. The merge issuing a chunk's loads first (c817f13c) changes no bits (the same 4,092 logit rows, shas and
  round counts) and takes the merge from 10.0 to 4.6 µs a call on GB10, attention 199.2 ms in all.
• Speed in this PR: none claimed, nothing calls it.
• Tests run: zig build test and zig build test-cuda-host (27/27 steps, 79/79 tests). zig fmt --check clean.
  lean_check.py: 538 findings, as on main.
• Not run, and why: Metal (unchanged); no model on main uses it.

Registration lines in build/cuda.zig, cuda/kernels.zig, cuda/root.zig and tests/cuda/main.zig sit where
#734, #736 and #740 add theirs; whichever lands second takes both.

Checklist

• Every token still goes through the lane rounds. No serial path, and nothing that needs drafts off.
• Drafted output equals "draft": false, a resumed prompt equals a fresh one, and concurrent equals solo. (No
  serving path changes.)
• No precision traded for speed. If the bits change, the description says which sums change and why.
• Prompt processing is no slower than the last release.
• The description says which platforms I ran, Metal M1 to M5 and CUDA, and which I could not.
• New tests fail before the change, pass after it, and skip cleanly without their dependency.
• Comments and docstrings are one line. No measurements or history in the source.
• No personal data, machine names, internal hosts or local paths. No AI attribution lines.