[jayhesselberth]

On sm_80+ an fp16x2 kernel scores two references per thread for every read inside the new cuda::fits_f16 bound (min(n,L)·m + 2|o| + |e| ≤ 2048, ~420 for tRNA); the int kernel takes the rest in the same batch. Output is byte-identical: gpu_parity runs every check on both kernels, including positive-mismatch and free-gap schemes, and both deliberately broken variants fail it. Rebased onto #11 (its narrow tail stripe stays on the int kernel; the half2 kernel still walks a full 16-row tail — porting RR ∈ {16, 8, 4} to it is the obvious follow-up).

A30, examples/gpu_score_probe.rs, M1 reads vs the 164-ref panel, int kernel as of #11 vs half2:

• 100–200 nt: local 2.58e11 → 5.28e11 cells/s (2.05×); semi-global 2.96e11 → 5.77e11 (1.95×)
• all of M1: local 5.32 → 2.61 s, semi-global 4.63 → 2.39 s (2.04× / 1.94×; against the pre-perf: non-saturating i16 adds in CPU kernels, narrow CUDA tail stripe #11 int kernel the spike was gated on, 2.15× / 2.04×)
• end to end, --device gpu, pre-perf: non-saturating i16 adds in CPU kernels, narrow CUDA tail stripe #11 main vs this branch: 6.4 → 4.3 s wall (md5 equal); 10.3 → 5.5 s with --strand both --secondary --mode semiglobal
• __launch_bounds__(256, 3) = 80 registers on evidence: 64 and 72 spill, 96 loses occupancy (table in benchmarks/README.md)

Checks after the rebase, all in srun allocations: fmt, clippy with and without --features gpu, workspace tests 130/130 (rna); gpu_parity 5/5 (288,480 pairs equal on both kernels), CLI GPU tests 5/5, cuda unit tests 2/2 (A30).

Closes #9

🤖 Generated with Claude Code