perf(cuda): gdn_scan_step_f32 kernel-internal rewrite — fuse decay+retrieve and write+read passes #173

Continuation of the non-GEMV decode-overhead investigation (issues #157/#159/#161/#164/#166/#168/#170, all merged to dev). Post-#170 profile: gdn-5-scan is the single largest non-GEMV decode bucket (~22.96ms / 4 profiled decode steps, 192 calls = 48 GDN layers x 4 steps), unchanged across #168/#170 because during decode (seqLen==1) gdn_scan_step_f32 is already exactly one kernel launch per (GDN layer, decode step) — no more launch-fusion win is available (see .docs/handoff.md Next Steps #1).

This issue targets a genuinely new angle for this specific kernel: a kernel-internal restructure of gdn_scan_step_f32 itself (native/kernels/gated_delta_net_scan.cu), never attempted before in this investigation (prior negative results are almost all on the separate pq2_0_gemv.cu kernel).

Analysis (SASS-first, per investigation discipline)

nvcc -Xptxas -v on the current kernel: 40 registers/thread, 0 spills, 1 barrier resource, blockDim=128 (4 warps/block). At 40 regs x 128 threads = 5120 regs/block, the SM's 65536-register budget floors to 12 blocks/SM = 48 warps/SM, exactly matching the sm_86 1536-thread/SM (48-warp) ceiling. Occupancy is already 100% — there is no occupancy headroom to gain here; any real win must reduce per-thread instruction/memory-op count instead.

cuobjdump --dump-sass confirms 2 BAR.SYNC instructions in the compiled kernel (matching its 2 source __syncthreads() calls).

Re-deriving the kernel's own dependency structure against both its header comment and the CPU oracle (DotLLM.Cpu.Kernels.GatedDeltaNetScan.Execute) row-by-row: every phase (decay, retrieve, write, read) has thread col touch only S[*, col] — its own column, exclusively, for the entire kernel body (the decay phase's grid-stride loop over the linearized [d_state, d_state] buffer reduces, since blockDim.x == d_state, to exactly row = i / d_state, colIndex = i % d_state == col). No other thread ever reads or writes column col. This means:

1. The second __syncthreads() (between write and read) is provably redundant — same-thread same-address load/store ordering is already guaranteed by the execution model, with zero cross-thread dependency on S after the first sync (which is only needed for k_shared/q_shared cross-thread visibility).
2. Decay writes S[row,col] and retrieve immediately re-reads it (separately, today) — since decay is a strictly order-independent per-element multiply, decaying S[row,col] and immediately consuming that same register value for the retrieve accumulation (same row-order 0..d_state-1 the CPU reference uses) is bit-identical to the current two-pass version, but eliminates retrieve's redundant reload.
3. Same argument for write followed by read: write's S[row,col] += k[row]*tmp[col] is a single order-independent per-element update; fusing it with the read-phase accumulation (again in the CPU's exact row order) is bit-identical but eliminates read's redundant reload.

Net effect: 4 full passes over the per-head [d_state,d_state] state matrix collapse to 2 fused passes, cutting LDG instruction count roughly in half (STG count unchanged — writes are unavoidable), and removing 1 of 2 __syncthreads() per kernel invocation, with proven bit-exact parity (verified by new xUnit test comparing GPU output to the CPU reference across multiple decode steps, since no such focused CUDA test existed before this issue).

Plan

• Rewrite gdn_scan_step_f32 with the two fused loops described above.
• Add tests/DotLLM.Tests.Unit/Cuda/CudaGdnScanStepF32Tests.cs: multi-step bit-exact comparison against DotLLM.Cpu.Kernels.GatedDeltaNetScan.Execute.
• Re-check ptxas -v register/spill count post-change.
• Measure real end-to-end dotnet run ... bench throughput (not just the DOTLLM_HYBRID_PROFILE=1 category profiler, which overstates real impact 3-10x per this investigation's established findings) — real Bonsai-27B GGUF, -p 64 -n 48 -r 8, A/B against a clean baseline build.
• Document result (positive or negative) in the kernel file's header comment, matching the existing documented-negative-results convention.