Skip to content

perf(cuda): narrow tail stripe in the half2 kernel - #14

Merged
jayhesselberth merged 1 commit into
mainfrom
perf-cuda-half2-tail
Oct 17, 2026
Merged

jayhesselberth merged 1 commit into
mainfrom
perf-cuda-half2-tail

Conversation

@jayhesselberth

Copy link
Copy Markdown
Member

Ports #11's narrow tail stripe to the half2 kernel: stripe2 is now a template on RR ∈ {16, 8, 4} and the last stripe takes the narrowest that covers the read's remaining rows (rows ≤ 4 → 4, ≤ 8 → 8), so a 130 nt read computes 132 rows instead of 144. The int kernel, MAX_REF_LEN_H2, fits_f16 and the launch bounds are untouched.

Measured (examples/gpu_score_probe.rs, one A30, M1 reads vs the 164-reference panel, batch 50,000, auto, arms interleaved, two reps each within 0.3%; fixed-length sets are M1 reads truncated to that length, 100,000 each):

reads tail rows (stripe) local main → tail × semi-global main → tail ×
100 nt 4 (4) 5.02e11 → 5.57e11 1.11 5.48e11 → 6.21e11 1.13
112 nt 16 (16) 5.61e11 → 5.62e11 1.00 6.12e11 → 6.30e11 1.03
120 nt 8 (8) 5.27e11 → 5.59e11 1.06 5.75e11 → 6.27e11 1.09
128 nt 16 (16) 5.61e11 → 5.63e11 1.00 6.13e11 → 6.31e11 1.03
130 nt 2 (4) 5.07e11 → 5.51e11 1.09 5.55e11 → 6.16e11 1.11
136 nt 8 (8) 5.30e11 → 5.60e11 1.06 5.80e11 → 6.28e11 1.08
144 nt 16 (16) 5.60e11 → 5.63e11 1.00 6.12e11 → 6.33e11 1.03
160 nt 16 (16) 5.60e11 → 5.62e11 1.00 6.13e11 → 6.32e11 1.03
M1, 100–200 nt — 5.28e11 → 5.49e11 1.04 5.77e11 → 6.15e11 1.07
M1, all — 5.25e11 → 5.48e11 (2.61 → 2.50 s) 1.04 5.74e11 → 6.13e11 (2.39 → 2.24 s) 1.07

Registers: __launch_bounds__(256, 3) unchanged, half2 kernel still 80 registers; local memory per thread 56 → 72 B (the int kernel holds 120 B at 64). New GpuScorer::{local_bytes, local_bytes_h2} report it and the probe prints it.

Parity: gpu_parity 5/5 on both kernels, both modes, the five schemes (including the positive-mismatch and free-gap ones only half2 masking survives), every read length 1–40 around the stripe; with the rows ≤ 8 arm deliberately sent to the 4-row stripe, 3 of its 5 tests fail (e.g. a 600 nt read scores 247 against the oracle's 251), so the test sees a wrong tail. CLI GPU tests 16/16 including device_gpu_output_is_byte_identical and pair_gpu_matches_cpu (A30); fmt, clippy with and without --features gpu, workspace tests and doctests on rna.

Docs: cuda/mod.rs, cuda/kernel.rs, benchmarks/README.md (new section with the table and reproduce lines), CHANGELOG.md (Unreleased), crates/eschalign/CLAUDE.md.

Closes #13

The last stripe takes the narrowest of 16, 8 or 4 rows that covers the
read, as the int kernel's has since #11: 1.09x local / 1.11x semi-global
at 130 nt on an A30, still 80 registers. GpuScorer reports local memory.
@jayhesselberth
jayhesselberth merged commit a5156a9 into main Oct 17, 2026
5 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

perf(cuda): narrow tail stripe in the half2 kernel

1 participant