Repository navigation
feat(qwen4): select QSA indexer blocks and their tail in native CUDA - #1178
Conversation
复核 + 一处 API 收紧在 agent 交付后按数据流与门禁纪律复核了一遍:
修改:把 gate 的 variant 从生产入口剥离(
|
|
This pull request has merge conflicts with main. Please rebase it so it can be reviewed and merged. |
Signed-off-by: bohaotu <bohaotu@pony.ai>
…entry The only public entry took the gate's mutation variant as a required argument, so a production caller had to spell a value that must always be 0 and could silently select a negative-control path. Split it the way the PLE block does: `qsa_indexer_select_into` fixes the production variant, and `qsa_indexer_select_variant_into` carries the knob for the parity gate. The kernel is unchanged. Also record the trap binaries' device side effect: on an RTX 5060 under CUDA 12.9.1 a deliberate `__trap()` left the whole GPU uninitializable for later processes, so they must run last or on a disposable device. Signed-off-by: bohaotu <bohaotu@pony.ai>
- Treat an empty launch as a no-op in the selector, matching the score stage and the other kernels' convention; the C entry still refuses one reached directly. - Count the selector's static shared state against the 48 KiB per-block ceiling so an exact-boundary launch cannot be refused by the driver. - Drop the dead int64 overflow guard: an int x int product always fits. - Correct the budget error text and the module/feature docs, which still described selection as a later stage. - Trap a negative `rows` entry in its own binary; the negative half of the row-position guard was untested. - Record in the bring-up doc that ties at the cut are not an edge case at long context, so an end-to-end accuracy gate must not require mask equality with the reference there. Signed-off-by: bohaotu <bohaotu@pony.ai>
One indexer, one golden gate: scoring and selection now live in `qwen4_qsa_indexer_parity`, sharing the fixture path, provenance check and tensor readers instead of duplicating them across two files. All five tests and every assertion are preserved; the select half still replays the fixture's recorded scores. The five trap binaries keep one trap per process, because `__trap()` leaves the context poisoned. Their bodies move into the shared `qwen4_qsa_trap` launcher, leaving each binary about eighteen lines like pegainfer-project#1160's, and the device-reset warning is recorded once in the launcher. Signed-off-by: bohaotu <bohaotu@pony.ai>
The five trap binaries each carried a case, but the cases are one line of setup. Keep one trap per process — `__trap()` poisons the context — and let one gate re-execute itself with `QSA_TRAP_CASE`, so the indexer has two test files: the golden parity gate and this trap gate. The child sets `PEGAINFER_REQUIRE_GPU=1` so a missing device fails instead of reporting a trap it never reached. Signed-off-by: bohaotu <bohaotu@pony.ai>
Signed-off-by: bohaotu <bohaotu@pony.ai>
d44a8f1 to
8fc98e4
Compare
FeathBow
left a comment
There was a problem hiding this comment.
The selection reads right against the reference's own mask, and the score gate moved into the shared file with its assertions unchanged. A few things before this leaves draft, details inline.
The description no longer matches the branch: the gate is now qwen4_qsa_indexer_parity, the traps run from qwen4_qsa_indexer_trap, the wrapper treats an empty launch as a no-op, and the PR also restructures the #1160 tests. It also says the device rerun after the API split is still pending, and the kernel changed after that. Could you rerun the gate and the trap binary at this head and update the numbers and commands?
| score_capacity.is_multiple_of(query_rows), | ||
| "QSA indexer score buffer {score_capacity} f32 does not divide into {query_rows} rows" | ||
| ); | ||
| let score_cols = score_capacity / query_rows; |
There was a problem hiding this comment.
The score stage takes score_cols explicitly and accepts a buffer longer than rows * score_cols; here the stride is inferred from the length. Handing this stage the score stage's buffer with spare capacity gives a larger stride, passes both checks, and selects from the wrong rows with no error. Could this take score_cols like the score stage?
| // opt-in. | ||
| constexpr int kMaxSharedBytes = 48 * 1024; | ||
|
|
||
| enum SelectVariant : int { |
There was a problem hiding this comment.
These variants stay in the production kernel and its C ABI, with six constants exported from ops, for something only the gate uses. The table they produce is good evidence, and temporary patches give the same evidence, which is how the #1160 review measured its gate. Could the knob go, with the table kept in the description?
| .output() | ||
| .unwrap_or_else(|err| panic!("spawn {case}: {err}")); | ||
| assert!( | ||
| output.status.success(), |
There was a problem hiding this comment.
Only the exit status is checked. If --exact trap_case stops matching (a rename, a harness flag), the child runs zero tests and exits 0, so every case passes without reaching a trap. Checking that the child reports one passed test would close that.
|
|
||
| __shared__ float warp_score[kWarps]; | ||
| __shared__ int warp_index[kWarps]; | ||
| for (int pick = 0; pick < keep; ++pick) { |
There was a problem hiding this comment.
Each pick is a full pass over the row's blocks plus two barriers: up to 512 passes over 65536 blocks per row at the checkpoint's 262144 positions, and blocks passes for every row shorter than the budget, where every block is kept anyway. Has this been timed at long context? A fast path for rows that keep all their blocks is cheap at least.
The selector inferred `score_cols` from `scores.len() / rows.len()`. The score stage allows a buffer with spare capacity, so handing the selector that buffer produced a larger stride, passed both host checks and read every row from the wrong offset. Take `score_cols` like the score stage does, check `scores.len() >= rows * score_cols`, and cover the spare- capacity case in the launcher tests. Signed-off-by: bohaotu <bohaotu@pony.ai>
The kernel, its C ABI, the ops surface and six exported constants carried a gate-only mutation parameter. Measure the five controls with temporary kernel patches instead, as the pegainfer-project#1160 gate did, and keep the production launch free of the knob. The golden, tie and structural assertions stay; the fixture still covers both selection paths. Signed-off-by: bohaotu <bohaotu@pony.ai>
Rows under the budget keep every complete block, so the mask is the whole visible prefix. Skip the masked argmax loop and its barriers for them; the parity gate already counts those rows and asserts the fixture covers both paths. Signed-off-by: bohaotu <bohaotu@pony.ai>
A filter that matches nothing (`--exact` after a rename) exits 0 with `0 passed`, so checking only the exit status let every case pass without reaching a trap. Require the child to report one passed test. Signed-off-by: bohaotu <bohaotu@pony.ai>
The masked-argmax selection cost one full row scan and two barriers per kept block: 86 ms for a single 256k-token row, 18 ms at 64k and 0.9 ms at the first truncating row. Replace it with an exact 8-bit radix select: four histogram passes find the keep-th sortable key, then the blocks equal to the threshold are filled in ascending block-index order (the canonical tie rule). The all-kept fast path stays. Measured on one L20, one row, budget 2048, nsys kernel medians at 256 threads: 20.8 us at 2052 tokens, 34.4 us at 8k, 67.0 us at 64k, 298 us at 256k, and 1.4 us for the all-kept fast path. The parity gate passes 5/5 and compute-sanitizer memcheck/racecheck are clean. Add the manual benchmark example used for those numbers. Signed-off-by: bohaotu <bohaotu@pony.ai>
Signed-off-by: bohaotu <bohaotu@pony.ai>
Clippy's unreadable_literal gate rejects 65535 and 262143 without separators; use separators so the manual benchmark stays clean under -D warnings. Signed-off-by: bohaotu <bohaotu@pony.ai>
Selector tuning: radix select instead of 512 argmax passesThe inline long-context question was right. The truncating path was 512 masked argmax passes, each a full row scan plus two barriers; at 256k tokens that was 86 ms for a single row. I replaced it with an exact 8-bit radix select:
One row, budget 2048 (before = idle host p50; after = nsys kernel median at 256 threads):
Thread sweep (nsys on a shared node): 128 threads is better at 2k/8k, 512 at 256k (187 vs 298 us); 256 is the balanced default. Correctness: the parity gate passes 5/5 at this head; selection is 133 rows, 129 bitwise, 4 up to a tie, 128 all-kept, 5 truncating. Benchmark: Environment caveat: the L20 image ships CUDA 12.9, below the repo's CUDA 13 floor, so this gate build used the FlashInfer tree's CCCL and a build-only guard for |
Parity + trap rerun at this head (L20)Ran both gates on one L20 (sm_89, CUDA 12.9 plus the FlashInfer tree's CCCL and the build-only
The gates ran on GPU 5 of a shared node with other tenants active. The PR stays draft; the description's Verification and Boundaries now reflect the trap run. |
FeathBow
left a comment
There was a problem hiding this comment.
All four points from last round are in: score_cols is explicit, the knob is gone from the kernel and the ABI, the trap parent checks for one passed test, and the radix select brings the 256k row down to 298 µs. The radix select reads right.
What is left is what the branch carries around the kernel, and we want that cleared out: comments and doc text that describe the argmax version or repeat each other, test scaffolding, a benchmark example, and a rebuild of the score stage's tests that this feature does not need. Details inline.
| return (bits & 0x80000000u) ? ~bits : (bits ^ 0x80000000u); | ||
| } | ||
|
|
||
| // One bit per block, set once a block has been picked. The bitmap is the whole |
There was a problem hiding this comment.
This describes the argmax version. The bitmap now only marks which threshold equal blocks get taken, so "exclusion state" and "picked" no longer fit. The same leftover sits in the shared memory ceiling comment ("exclusion bitmap") and the description.
| // block tied with the last kept one is an equally valid pick; this kernel's | ||
| // tie rule is the canonical one the gate holds the fixture to. | ||
| // | ||
| // A non-finite score inside `[0, blocks)` is an upstream contract violation and |
There was a problem hiding this comment.
The selection contract is spread over four docs (this header, the C entry, the FFI declaration and qsa_indexer_select_into), and this header and the wrapper doc each give the full version: grid contract, tie rule, the torch.topk note and the trap. The NaN rationale alone appears here, again in the kernel at the finite check, and again in the wrapper doc. Much of it argues why the design is right (why trap instead of skip, why the kernel repeats the host checks, why the host mirrors MAX_KEPT_BLOCKS), and the header also carries derived numbers (319392 blocks, 1.28M tokens). Could this header be the one place for the semantics, tie rule and trap, in a few lines, with the wrapper doc keeping only the Rust side (stride, spare capacity, empty launch) and the rationale moving to the description?
| // | ||
| // The compiled geometry is kRatio = 4 and one byte per mask cell; a caller | ||
| // whose buffers are laid out differently cannot be detected here. An empty | ||
| // launch (no rows, or no tokens) is refused rather than treated as a no-op: a |
There was a problem hiding this comment.
This says an empty launch is refused because asking for nothing is a scheduling bug, but the only caller, qsa_indexer_select_into, returns early on exactly that case and treats it as nothing to do. The check can stay; the justification contradicts the wrapper and can go.
| score_cols: i32, | ||
| } | ||
|
|
||
| /// The launcher's pre-launch contract. Split from the launch so the guards can |
There was a problem hiding this comment.
select_plan is split out so its guards can be unit tested, and the five tests check one line ensure!s. select_plan_refuses_empty_extents checks a branch the public entry never reaches, since it returns early on an empty launch. The score stage above has no such split. Could select_plan fold back into qsa_indexer_select_into and these tests go?
| //! prefix-evicted visible list is out of contract and has no case here. | ||
| //! | ||
| //! The fixture covers both selection paths: rows that keep every block take | ||
| //! the all-kept fast path, and the truncating rows run the masked argmax |
There was a problem hiding this comment.
This still says masked argmax, and the module doc above carries this fixture's measured counts (133 rows, 129 bitwise and so on), which go stale with any fixture change and belong in the description.
More broadly, the score half here is the merged score parity file moved into mod score and reindented, with its assertions unchanged, so 750 lines are deleted and readded for code this PR doesn't change. Could the selection test sit beside the existing score test in its file, without moving the score half?
| //! | ||
| //! `__trap()` leaves the CUDA context in a sticky error state, so a process can | ||
| //! fire one trap and no more. This binary runs each case by re-executing itself | ||
| //! with `QSA_TRAP_CASE` set; the parent only spawns and checks the exit status. |
There was a problem hiding this comment.
This replaces the score stage's merged trap binaries with a harness that reruns itself per case, and adds three selection traps. The header still says the parent only checks the exit status, and it carries a device model and an incident report. The selection traps cover a range check and a finite check that are a few lines each, they have to fault the device, and nothing runs them outside a disposable card. Could the selection traps go and the score stage's trap files stay as they are on main?
| @@ -0,0 +1,179 @@ | |||
| //! Manual timing for the QSA indexer selection kernel. | |||
There was a problem hiding this comment.
This reproduces the description's timing table and checks only the mask's set bit count. The numbers already live in the description, so the example and its [[example]] entry in Cargo.toml can go, along with the bench command in bring-up.md.
|
|
||
| ## The QSA indexer's selection half | ||
|
|
||
| `csrc/qwen4/qwen4_qsa_indexer_select.cu` (`qsa_indexer_select_into`) turns those scores into the token-level mask the sparse attention consumes — one `u8` per (row, token), row stride `tokens`, which is the same contract the attention unit of #1105 reads. A row at `p` keeps the `min(budget/4, blocks)` highest block scores (descending score, ascending block index — the canonical tie order), sets their four tokens each, and sets the tail `4*blocks ..= p` unconditionally, so causality is inherent rather than an extra condition. The incomplete tail is exactly why the scoring pass never scores it. One CTA per row: a row under the budget keeps every block and writes the whole visible prefix in one pass; otherwise an exact 8-bit radix select over the row's sortable f32 keys (four passes) finds the `keep`-th score, and the blocks equal to that threshold are filled in ascending block-index order — the canonical tie rule — with a one-bit-per-block bitmap. The target is the same portable subset (no TMA, no cluster, no wgmma), and the bitmap plus the kernel's static state (per-warp radix histograms, threshold and equality prefix) fit under the 48 KiB per-block ceiling, which covers 1.28M tokens, past this checkpoint's 262144 positions. It inherits the scoring pass's token-0 grid contract: a row at `p` sees exactly `[0, p]`, so a prefix-evicted visible list needs a shifted key array on both stages and is not supported here. The wrapper treats an empty launch as a no-op (the scoring stage does the same) and refuses a budget that is not a positive multiple of four or keeps more than 512 blocks, a score row narrower than the row ceiling and a short buffer; a non-finite score inside a row's prefix or an out-of-range `rows` entry traps on the device — zeroing or skipping a poisoned block would hand the attention softmax a shorter visible set with no trace, and an all-masked row is undefined there. |
There was a problem hiding this comment.
This paragraph is a fifth copy of the kernel contract: mask layout, tie rule, the radix algorithm, the shared memory bound, the grid contract, the wrapper's refusals and the NaN rationale. The timing paragraph below repeats the description: the before and after table, the thread sweep, the build environment and the mutation table. The doc could keep what the code can't say: the radix select result at 256k, the token 0 grid limitation, and the note that an accuracy gate must not require mask equality at a tied cut.
Review round appliedAll eight inline points are in at
Verification at this head ( |
The score stage's test files, parity gate and trap binaries are restored byte-for-byte from main; selection tests sit beside the score test in `qwen4_qsa_indexer_score_parity.rs` without moving the score half. The select kernel header is the one place for the contract; the C entry and the Rust wrapper keep only what each side owns. The `SelectPlan` split and its five device-free tests fold back into the wrapper, the benchmark example and its Cargo entry go, and the bring-up doc keeps only what the code cannot say: the radix select's 256k result, the token-0 grid, and the tied-cut rule. Signed-off-by: bohaotu <bohaotu@pony.ai>
Signed-off-by: bohaotu <bohaotu@pony.ai>
e477e57 to
fadce50
Compare
…ted docs Signed-off-by: Feathbow <feathbow@gmail.com>
Commit attribution noticeThis pull request contains commits attributed to an account other than the PR author @sysuTobo without a verified signature from that account:
If you are mentioned above and did not take part in this pull request, please say so here. Attributing a commit to someone else normally requires that person to be involved; maintainers should treat unacknowledged attribution as a review blocker. This check never fails the build. |
|
We pushed one commit on top that trims what the last review pointed at. Selection behaviour is unchanged.
Checked on sm_89 with the branch merged into current main: fmt, clippy with The description still mentions the device-free guard test. |
Summary
QSA unit three of #1105, after the scoring unit (#1160): turn the indexer's per-(row, block) scores into the token-level visibility mask, in native CUDA.
csrc/qwen4/qwen4_qsa_indexer_select.cuadds one CTA per query row that writesu8 [query_rows, tokens], row stridetokens— the shape #1105's sparse attention consumes. Wrapperqsa_indexer_select_intoinsrc/ops/qwen4.rs, FFI insrc/ffi/qwen4.rs, selection gate in theselectmodule oftests/qwen4_qsa_indexer_score_parity.rs, and a bring-up doc update. The scoring stage's parity file and its two trap binaries are byte-for-byte as on main; only the selection module is added beside them.Selection is an exact 8-bit radix select over the row's sortable f32 keys: four histogram passes find the
keep-th score, the blocks above it are kept, and the threshold-equal blocks fill the rest in ascending block-index order (the canonical tie rule). A row under the budget keeps every block and writes the visible prefix in one pass. This replaces 512 masked argmax passes — 0.9 ms at the first truncating row, 18 ms at 64k and 86 ms at 256k — with 20.8 / 67 / 298 µs on one L20.Exact semantics
For a query row at position
p(token 0 is the first visible token — the scoring pass's grid contract):blocks = (p + 1) / 4complete micro-blocks,tail = (p + 1) % 4tokens at4*blocks ..= p;keep = min(budget / 4, blocks), withbudgeta positive multiple of 4 keeping at most 512 blocks;keephighest block scores are kept under the canonical total order descending score, ascending block index;p.score_colsis an explicit argument, like the score stage:scores.len()may exceedquery_rows * score_cols(spare capacity is allowed) and the row stride staysscore_cols.A row under the budget keeps every complete block and one pass writes its visible prefix. Otherwise the selector maps each finite score to a monotonic u32 key (
-0.0folded to+0.0, so both zeros tie), runs four 8-bit histogram passes to find thekeep-th key, keeps every key above it, and fills the remaining slots from the equal keys in ascending block-index order. A non-finite score inside a row's block prefix traps rather than being skipped: the reference score is a relu sum, so a poison means the indexer input is already broken, and silently dropping that block would hand the attention softmax a shorter visible set with no trace (an all-masked row is undefined there). The wrapper treats an empty launch (rows.len() == 0ortokens == 0) as a no-op, matching the score stage; the C entry still refuses one reached directly. The wrapper refuses a malformed budget and a short buffer; the kernel repeats the budget and score-row-width checks and traps on out-of-rangerowsentries.No gate knob
The production kernel, its C ABI and the
opssurface carry no mutation variant. The five mutation controls in the table below were measured by temporarily patching the kernel, as the #1160 review did; they are not shipped code.What the gate measures
Selection is comparison-only over the fixture's recorded scores, so there is no tolerance: strict-cut rows must match the reference mask bitwise. A block tied with the last kept one is an equally valid
torch.topkpick, so a tie row is held to: both masks structurally valid (causal prefix, complete 4-token blocks only plus the exact tail, exactlymin(K, n)*4 + tailbits set), every selected block scoring at least every excluded one, and the same number of tied-at-cut picks. Tie rows are counted and reported, never skipped.Measured on one L20 (sm_89, CUDA 12.9 plus the flashinfer tree's CCCL, 256 threads): 5/5 tests pass; 133 rows, 129 compared bitwise, 4 compared up to a tie, 0 mismatches; the 4 tie rows are case 6 at cut score 0.0 with 31–43 tied blocks; the fixture covers 128 all-kept and 5 truncating rows.
The five mutation controls, measured by temporary patches:
Performance, one row, budget 2048 (old = idle host p50; new =
nsyskernel median at 256 threads):Thread sweep on the shared node (
nsyskernel medians): 128 threads is faster at 2k/8k (idle host p50 17.9 / 26.9 µs against 20.8 / 34.4 µs at 256), 512 is faster at 256k (187 against 298 µs); 256 is the balanced default. These numbers were measured with a temporary benchmark binary that does not ship.Device-free coverage: the
selectmodule keeps one test for the C entry's guards (null pointers, zero/negative extents, malformed budget, overly narrow score row, oversized bitmap), which answerCUDA_ERROR_INVALID_VALUEbefore any pointer is formed. The Rust-side guards are inline inqsa_indexer_select_into(bad budget, short buffers, spare capacity, overflowing products).Verification
cargo fmt --all -- --check,cargo check -p pegainfer-kernels --features qwen4 --tests, andcargo clippy -p pegainfer-kernels --features qwen4 --tests -- -D warningspass at this head. The build used CUDA 12.9 plus the flashinfer tree's CCCL and a build-only guard for the CUDA-13-onlycudaErrorLossyQuerycase; the production CUDA 13.2 build is the CI/dev image and was not re-run here.compute-sanitizermemcheck on the 256k case: 0 errors; racecheck on the 2k+ case: 0 hazards. (Measured at the previous head; this head renames the bitmap's helpers and trims comments, and the selector logic is unchanged.)nsyskernel medians: 25 launches of the 256k case, median 298 µs (min 297.5, max 298.8); the shared node's other tenants are why the thread sweep uses nsys medians.ncuis unavailable on the shared pod:ERR_NVGPUCTRPERM(no GPU performance-counter permission).Boundaries not gated
qwen4feature, so CI neither builds nor runs any of this; the gate is manual (PEGAINFER_REQUIRE_GPU=1).rows, non-finite score) are not exercised by a test in this PR; the score stage's trap binaries are unchanged from main and were not re-run here.Refs #1103 #1105.