[Bizuayeu]

What this changes

What it buys: a ReleaseSafe build of the contract's runner reaches its SUMMARY, with the default build's results,
instead of stopping three times. Only the contract's test adapters change, so prompt and decode speed and the served
bytes cannot move.

zig build kernel-contract builds its runner at the build's default, fast. Built ReleaseSafe instead (as one does to
check a runner), the CUDA contract stops three times before its SUMMARY, each in a test adapter's host code:

• Fexp.drop() in zig/tests/cuda/contract_flashnext_fexp.zig freed its expert list ([]usize) and its gate and up
  scale rows ([]f32) through std.mem.sliceAsBytes, at byte alignment, so the repeat gate's second prepare of
  "flashnext factored experts down 64 pairs" panics: free of [addr: ef747402f070, len: 64 (0x40) align: 1] mismatches allocation of [addr: ef747402f070, len: 64 (0x40) align: 8] (contract_flashnext_fexp.zig:108 in
  prepare, from gates.zig:172 in repeat). drop() now frees each slice at its own type; freeing an empty slice
  does nothing, so the length check goes too.
• Rows.stress() in zig/tests/cuda/contract_window_rows.zig copied the first cache row onto every row starting
  with the first, so its first @memcpy had the same source and destination: @memcpy arguments alias
  (contract_window_rows.zig:249 in stress, from gates.zig:284 in sums, the window_rows decode case). The loop
  now starts at the second row; every cache holds at least one row (crow is the ring, the prompt's keys and rows,
  or streams × ring).
• Chunks.attend() in zig/tests/cuda/contract_glm_attn.zig (mine, from glm/cuda: KV pages from the core pool, kept states sized by their position (#1142) #1163) keeps a row's scores, their bounds
  and its probabilities in one scratch of 3n and looped the probabilities over the whole scratch: for (p, s) with
  p n long and s 3n, for loop over objects with non-equal lengths (the GLM chunk attention's nonfinite gate, on
  a page table out of order). It now loops over the scores' first n, s[0..n], which is what the fast build happened
  to read.

None changes what a gate computes: the first frees scratch at the end of a gate, the second skips a copy of a row
onto itself, the third reads the n scores the probabilities come from.

Receipt

• Environment: TensorFold fa735852d (main 1b4f2f979 and the three commits) against main 1b4f2f979; GB10 (sm_121,
  128 GB unified memory); driver 580.178.04; nvcc 13.3.1; tf-cuda-test built with zig build compile-cuda-tests -Doptimize=ReleaseSafe -Dnvcc=..., run as tf-cuda-test kernel-contract.
• Before, main at ReleaseSafe: the run stopped at the fexp panic after 722 PASS and 76 FAIL lines (74 planted
  controls going red, and Qwen 3.5's two decode attention nonfinite lines the run marks KNOWN, CUDA kernel contract: the Qwen 3.5 decode attention cases launch tf_q35_attn_partial and tf_q35_attn_partial8 at 192 threads, past their __launch_bounds__(128), so every gate fails with CUDA_ERROR_INVALID_VALUE #1267). With the fexp
  fix alone it stopped at the window_rows panic after 744 PASS, the same 76 FAIL lines and no SUMMARY; with both,
  at the GLM attention panic after 1,286 PASS.
• After, all three fixes at ReleaseSafe: the run reaches its SUMMARY (5,821 s): SUMMARY 340 kernels: 340 pass gates 1-4, 0 fail, 5 known findings; 94 controls: 94 red as planted, 0 missed; 231 limit problems. Its 293 FAIL lines are,
  numbers aside, the same lines as a default-build run's on glm cuda: a decode step's select gives each row a 1,024-thread block #1424's branch, whose SUMMARY differs only by that
  branch's one added kernel, and both runs exit 1 on the same 231 limit problems (contract_tests.zig:594).
• Exactness, decode and prompt speed: n/a, tests' host code only; no engine code changes.
• Tests run: zig build compile-cuda-tests 10 of 10 steps, zig fmt --check, the lean check (no new problems
  against main).
• Not run, and why: Metal and HIP (these files are CUDA's contract).

Checklist

• Every token still goes through the lane rounds. No serial path, and nothing that needs drafts off. (No engine
  code changes.)
• Drafted output equals "draft": false, a resumed prompt equals a fresh one, and concurrent equals solo. (No
  engine code changes.)
• No precision traded for speed. If the bits change, the description says which sums change and why. (No
  arithmetic changes.)
• A new or changed GPU kernel is registered with the kernel contract (docs/kernel-contract.md); each backend's
  SUMMARY line is in the receipt. (No kernel changes; the CUDA SUMMARY line is in the receipt.)
• Prompt processing is no slower than the last release. (Not reached.)
• 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. (No new test; the
  repeat and sums gates that stopped before run through after.)
• 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.

[Bizuayeu]

Updated the description: it now opens with what the change buys (a ReleaseSafe runner reaches its SUMMARY; prompt and decode speed cannot move). No code change.

[ashhart]

Thanks, Bizuayeu. All three are real ReleaseSafe faults in the test adapters, and each fix is the smallest one that
works.

• Fexp.drop() frees the expert list and both scale rows at their own types. It still resets them to empty slices, so
  a repeated cleanup stays safe (zig/tests/cuda/contract_flashnext_fexp.zig:82-84).
• Rows.stress() copies the first cache row only into the rows after it, so no @memcpy copies a row onto itself
  (zig/tests/cuda/contract_window_rows.zig:249).
• Chunks.attend() computes probabilities from only the first n scores (zig/tests/cuda/contract_glm_attn.zig:198).
  It reads the same n scores in the same order, with no change to tolerance or reduction order.

We reproduced each panic on the host from the adapters' own code: the misaligned free, @memcpy arguments alias, and
the unequal-length for. Each one passes after your change. fmt, lean (no new findings) and git diff --check are
clean. The CUDA runner cross-compiles at ReleaseSafe, and the host contract steps pass 115/115. The branch applies
cleanly to current main.

The last check is ours to run. We're running the ReleaseSafe kernel contract
(zig build kernel-contract -Doptimize=ReleaseSafe) on a GB10, for main and for this branch, with the same nvcc. It
will land if this branch reaches SUMMARY with the same planted controls and known findings as main, and no new gate
or limit problems. You don't need to do anything more. Thank you!

[ashhart]

Landed on main in 20d1deb, with your three commits kept as they were (5d108e5, 733957a, 2b7eefc). Thank you! I ran the CUDA kernel contract in a ReleaseSafe build on a GB10 before landing: main's build stops at these adapters, and yours runs through to the end.

[ashhart]

This is in TensorFold 1.1.8, released today. Thank you, @Bizuayeu!