Repository navigation
perf(cuda): narrow tail stripe in the half2 kernel - #14
Merged
Merged
Conversation
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.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Ports #11's narrow tail stripe to the half2 kernel:
stripe2is now a template onRR ∈ {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_f16and 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):main→ tailmain→ tailRegisters:
__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). NewGpuScorer::{local_bytes, local_bytes_h2}report it and the probe prints it.Parity:
gpu_parity5/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 therows ≤ 8arm 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 includingdevice_gpu_output_is_byte_identicalandpair_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