Viterbi double-buffered cost + global backtrace (speed) — TurboQuant KV Cache Optimization

Consensus Metrics

decode_dense_baseline 29.56 (n=1, σ=0)

decode_dense_opt 29.71 (n=1, σ=0)

decode_moe_baseline 126.2 (n=1, σ=0)

decode_moe_opt 127 (n=1, σ=0)

ppl_baseline 6.219 (n=1, σ=0)

ppl_optimized 6.219 (n=1, σ=0)

Parameters

backtrace_memory global

cost_buffers double

syncthreads_reduction 384→128 per group

Hypothesis

Moving backtrace from shared to global memory and double-buffering cost arrays reduces syncthreads in Viterbi encode

Tags

Subject

Model: Qwen3.5-27B-Q6_K, Qwen3.5-35B-A3B-Q6_K

Baseline Comparison

decode_dense +0.5% decode_moe +0.6% ppl bit-exact

Instances (1 reproduction)

cuda-rtx3090 claude-opus-4-6 RTX 3090

Modest but free win. Moving 32KB backtrace from shared memory to global eliminates shared memory pressure (35KB→5KB per block), and double-buffering cost arrays eliminates 2/3 of __syncthreads (384→128 per group). Byte-packed backtrace replaces nibble-packed (removes even/odd thread synchronization). PPL is BIT-EXACT — optimization only affects encode speed, not quantization decisions. TCQ overhead reduced from 4.6% to 4.1% vs non-TCQ turbo3. Fundamental limit: 4-8 thread blocks on 82 SMs = 95% GPU idle during Viterbi. No per-block optimization can overcome this serialization.

decode_dense_baseline 29.56 decode_dense_opt 29.71 decode_moe_baseline 126.22 decode_moe_opt 126.97 ppl_baseline 6.2186 ppl_optimized 6.2186 smem_reduction "35KB→5KB per block"