[sysuTobo]

Required human contributor answers

1. What outcome are you trying to achieve? Lower long-prompt attention latency in the native Qwen3.5 CUDA backend, proven on the two GPUs we can measure (an L20 and an sm_120 workstation), with outputs that do not change.
2. What is the highest-leverage variable for that outcome? Occupancy and K/V reuse of the causal attention kernel. The kernel is capped at one block per SM by 67,584 bytes of dynamic shared memory, so the query tile alone decides warps per SM (4 or 8) and how many query rows share one K/V tile (64 or 128).
3. How will this work move that variable rather than create busywork? It changes only that tile: a 128-query, 256-thread instantiation with the same 32-key tiles, online-softmax recurrence and epilogue, so outputs stay byte-identical. It is enabled only on the two measured capabilities and for windows of at least 640 rows; the 64-query kernel, the shared-prefix kernel and every other architecture are untouched.
4. Who is the customer or user, and what do they actually need? The native workers that load libqwen3_5_cuda.so on L20-class (sm_89) hosts, plus sm_120 development workstations. They need faster long-prompt prefill without any change to decisions or probabilities.

After answering the four questions: is a PR needed at all to achieve this outcome? Yes. The change is measured, bit-identical and scoped to two capabilities; a PR is how it gets reviewed and qualified.

Bug fix? N/A. This is a performance change with byte-identical output.

Purpose

Long causal attention windows run four warps per SM because the block's 67,584 bytes of shared memory hold one block. A 128-query, 256-thread block with the same 32-key tiles raises that to eight warps per SM and reads each K/V tile once for 128 queries instead of 64, at 101,376 bytes (the sm_89/sm_120 opt-in limit), still one block per SM.

Each query row keeps the same key-tile walk, softmax recurrence and epilogue, so every measured output is byte-identical to the 64-query block; only which block computes a row changes.

Scope:

• Full pass (causal or bidirectional) and 128-aligned causal prefix windows of at least 640 query rows on sm_89 and sm_120.
• cs1_attention_gated_shared and every other architecture keep 64 queries. cs1_attention_gated_prefix uses the wider tile only when q_base % 128 == 0. The bidirectional entry point shares launch<> and widens the same way; it always has q_base = 0, and its GPU-gated reference test covers 700 and 2,048 rows.
• H200/sm_90 and end-to-end serving are unverified and stay on 64 queries.

Test Plan

• Frozen protocol, controls, budget and stop condition: benchmarks/attention-wide/plan.json.
• Kernel arms: both tiles in one process, two trials per device.
• Public entry point (cs1_attention_gated): base and candidate libraries swapped between arms, two trials per device.
• Parity: byte comparison of both arms at every length, and prefix windows at q_base 64/128/1,024 against the full pass.
• compute-sanitizer --tool memcheck at 640, 1,024, 4,096 and 8,192 rows.
• Dispatch: a CPU-stub run that alternates the current device (sm_89, sm_90, sm_120, then back) checks that the wide tile follows the current device and that the shared-memory opt-in is applied per device, plus an interleaved pre/post A/B of the public entry point to confirm the per-dispatch device query is not measurable.
• cargo test --release -p omni-qwen3-5-native --test kernels -- --ignored --test-threads=1 attention against the candidate library.

Test Result

Base 5fa5179d57f9267bb0d20162004d8ce26352aa9d plus this branch. Public cs1_attention_gated, old -> new, medians of two trials:

tokens	L20 (sm_89, 92 SMs)	RTX 5060 Laptop (sm_120, 26 SMs)
512	0.1300 -> 0.1296 ms (-0.3%)	0.7720 -> 0.7720 ms (0.0%)
640	0.1747 -> 0.1698 (-2.8%)	1.0478 -> 0.8279 (-21.0%)
1,024	0.3323 -> 0.3144 (-5.4%)	2.2597 -> 1.8092 (-19.9%)
2,048	0.9859 -> 0.8998 (-8.7%)	6.9033 -> 5.5222 (-20.0%)
4,096	3.4010 -> 2.9960 (-11.9%)	23.9185 -> 19.4021 (-18.9%)
8,192	12.0261 -> 10.4388 (-13.2%)	87.1708 -> 71.5387 (-17.9%)

Kernel arms alone: L20 -2.2% at 640 rows to -13.8% at 8,192, and +4.9% at 512 rows, which is why the tile is gated at 640; sm_120 -13.0% to -22.4% from 512 rows up. 512 rows is unchanged in the shipped dispatch because it stays on the 64-query block.

Fidelity: byte-identical outputs at every measured length in both trials on both devices; prefix windows at q_base 64/128/1,024 match the full pass; sanitizer reports 0 errors at all four lengths on both devices; the seven attention kernel tests pass, including windowed_attention_matches_full_pass_bit_for_bit and the float64 reference at 700 and 2,048 rows. The full ignored suite is 26 passed; the one remaining case cannot run here because QWEN3_5_CHECKPOINT is not set.

Demo / evidence

• Record and protocol: benchmarks/attention-wide/README.md
• Per-trial medians and delta ranges, parity and sanitizer outcomes: results.json
• Kernel change: src/backends/cuda/qwen3_5/attention.cu

The per-batch records from both runs are in commit a3bb8c4 on this branch; this tree keeps the frozen protocol and the per-trial results. plan.json fixes what the local measurement programs implemented. No HTTP, Public231 C1 or whole-model measurement is included, and nothing here is an end-to-end or task-accuracy claim.

Self-review

• RSI phase: kernel.
• Primary change: a second kernel instantiation with a 128-query, 256-thread block plus a launch-time choice of tile. The 64-query kernel is unchanged and still serves short windows, unaligned prefix windows, the shared-prefix path and unmeasured capabilities.
• Device ownership: the capability check and the shared-memory opt-in are per device and re-evaluated when a thread's current device changes, so cs1_set_device and mixed-device hosts do not inherit another device's decision.
• Alternatives considered: two blocks per SM via a swizzled layout and 16-key tiles (changes the softmax reduction order, not taken); 16-key tiles with the padded layout (103,424 bytes for two blocks, over the 102,400-byte SM budget); a 192-query tile (135,168 bytes, over the opt-in limit); lowering registers to fit a second block (shared memory, not registers, is the limit); enabling the wider tile at every length (512 rows is 4.9% slower on L20).
• Changed size: 220 added and 9 removed lines across two source files and the record, 48 of them in attention.cu.
• Validation by changed area: CUDA backend build, sanitizer, byte parity, the attention kernel tests above; the Rust crates are unchanged.
• Still unverified: H200/sm_90, the shared-prefix kernel, HTTP/Public231 C1, and whole-model parity on the changed tile.

[hsliuustc0106]

PR #202 first-pass review — "cuda: widen the attention query tile on sm_89 and sm_120"

Date: 2026-11-01. Reviewer pass: 1 (first review). VERDICT: COMMENTED (DRAFT PR — informational findings only, no P0/P1 found).

Metadata

• Head 4e69f6ca9 (4 commits), base main @ 5fa5179d5; merge-base verified 5fa5179d57f9... = main tip (freshly based).
• Author sysuTobo; OPEN, isDraft: true, no review decision. Author-reported "keeping this in draft until the owner reviews" — consistent.
• Diff: 5 files, +220/−9 — src/backends/cuda/qwen3_5/attention.cu (+48 net), backend README (+11), benchmarks/attention-wide/{README.md 72, plan.json 29, results.json 60}.

Observed CI — READINESS GAP (prominent)

Zero CI checks have executed anywhere for this PR.

• ThinkFlowLab/system1-omni: for head sha 4e69f6ca9..., actions/runs?head_sha= returns 2 runs — "CI" and "Docs", both conclusion: action_required (awaiting maintainer workflow approval; the author has no prior approvals on record). Same state for all four commits (8 branch-matched runs total, all action_required).
• Fork sysuTobo/system1-omni: 0 runs for the head sha and for the branch.
• Consequence: the cuda-compile job — the only CI tier that would compile flash_wide_kernel at all (this host has no nvcc/cargo; I could not compile locally either) — has never run on this code. Compile verification is currently author-reported only.

Source inspected

• Full PR body (saved at /tmp/pr-review-pass119/pr202-body.md), all 4 commit diffs, final attention.cu (whole file), plan.json, results.json, both READMEs, tests/qwen3_5/kernels.rs (unchanged from main), .github/workflows/ci.yml + docs.yml on main, benchmarks/gdn record layout (precedent), .agents/skills/self-review/SKILL.md.

Commands run

gh pr view 202, git log/show/diff/ls-tree/merge-base/grep against origin/pr/202/origin/main, gh api repos/ThinkFlowLab/system1-omni/actions/runs?head_sha=... and repos/sysuTobo/system1-omni/..., worktree at /tmp/pr-review-pass119/wt202 (removed after): PYTHONPATH=. /Users/hsliu/miniconda3/bin/python -m pytest tests/benchmarks -q → 133 passed (28s); /Users/hsliu/miniconda3/bin/python -m mkdocs build --strict → exit 0. No cargo/nvcc on this host.

Commit-by-commit hygiene

1. a3bb8c44 the change: +4,434 lines — kernel + record including raw results_*.jsonl (846×2), lib_*.jsonl, parity/sanitizer logs, three harness .cu programs, summarize.py, pr-description.md.
2. 7e5b1249 removes the three harness programs, keeps summarize.py.
3. 18049cb2 removes raw batches, logs, pr-description.md, summarize.py; hoists plan.json, renames summary.json → results.json.
4. 4e69f6ca final trim to 72-line README / 29-line plan / 60-line results + comment tightening in .cu.

Final tree is clean: benchmarks/attention-wide/ contains exactly README.md, plan.json, results.json; a name-level search for attn_record|attn_lib_record|attn_prefix_check|summarize|pr-description|*_l20|*_sm120|parity_|sanitizer_ over the whole head tree finds nothing. No orphaned measurement programs or one-off run output at head. (The 4.4k-line artifacts remain in history — normal for iterative cleanup, noted only.)

Kernel change — verification of the core claims

Same recurrence/epilogue is unchanged code, not rewritten. flash_body gains template params <..., int QROWS = BM, int NTHREADS = THREADS>; the diff substitutes QROWS/NTHREADS for BM/THREADS in exactly five places (qs/ks smem layout, q0 stride, Q-load loop, kv_end, K/V-load loops). The online-softmax block (mx → shfl_xor 1,2 → base/alpha → exp2f → l → o rescale), the P*V mma path and the epilogue (inv, gate rounding, pack_bf16 stores) are untouched. flash_wide_kernel only calls flash_body<Gated, false, Bidirectional, WIDE_BM=128, WIDE_THREADS=256>; flash_kernel keeps the defaults, so the 64-query instantiation is the identical code.

Byte-identity is structural, and the argument checks out beyond "same recurrence":

• Per row: row0 = q0 + warp*16; both 64 and 128 are multiples of 16, so for a row p, row0 = p − (p mod 16) is identical in both layouts — hence the per-warp active cutoff (k0 <= row0 + 15) and the exact sequence of key tiles that contribute to a row are the same. The extra tiles a wide block walks are active=false for that warp or fully masked; masked lanes contribute exp2f(−inf − base) = +0, +0 * v = +0, and acc + +0 is exact (alpha = exp2f(0) = 1.0 exactly on rescale). So the FP32 operation order per row is identical → same bits, structurally, not merely tested.
• Tested: tests/qwen3_5/kernels.rs::windowed_attention_matches_full_pass_bit_for_bit (unchanged from main) — seeded inputs, element-wise exact equality (assert_eq!(bad, 0)), no tolerance; on sm_89/sm_120 it cross-compares wide-tile full passes (t=712, 2048 ≥ 640) against 64-tile windows (q_base 64/896 give rows < 640). flash_attention_matches_float64_reference (t=700, 2048) uses a 1.6e-2 relative tolerance but is a separate accuracy check, not the basis of the identity claim. All are #[ignore = "needs a GPU and CUA_S1_CUDA_LIB"] → GPU-gated, not wired into CI (the CI rust job is GPU-less and runs without --ignored).

Dispatch/guard (launch<>): pre-validation unchanged (q_base % 64, q_base <= T, Bidirectional && q_base != 0 → error, etc.). Wide branch requires all three: wide_tile_supported() AND q_base % WIDE_BM == 0 AND rows >= 640; any miss falls through to the existing 64-query code path below it. Capability check uses cudaGetDevice + cudaGetDeviceProperties → (major==8 && minor==9) || (major==12 && minor==0) — sm_89/sm_120 exactly. Tail block (rows % 128 ≠ 0): Q rows ≥ T are zero-filled via the cp_async16 predicate and the epilogue skips row >= T. cs1_attention_gated_shared (shared-prefix kernel) has its own launch path, untouched.

Shared-memory bound: SMEM_BYTES = (64+2·32)·264·2 = 67,584 ✓; WIDE_SMEM_BYTES = (128+2·32)·264·2 = 101,376 = exactly 99 KiB, the sm_89/sm_120 opt-in per-block max (so cudaFuncSetAttribute at exactly the limit is legal), and one block/SM against the 102,400-byte per-SM budget → 8 warps/SM as claimed. Launch-failure path: attribute-set failure returns the cached cudaError_t; launch error → return cudaGetLastError() (clears and propagates) — no silent fallback to 64 queries, matching the existing pattern.

ABI: zero ops.h/Rust/exported-symbol changes; all extern "C" signatures unchanged; flash_wide_kernel lives in the anonymous namespace. Confirmed by diffstat (2 docs + 3 record files + 1 .cu).

Body vs record (cross-checked, all consistent)

• All 12 public-entry numbers in the body table match results.json library arrays to the shown precision (e.g. L20 8192 12.0261→10.4388 −13.199%; sm_120 640 1.0478→0.8279 −20.988%). Kernel-arm claims (+4.913% at 512 on L20; −13.008…−22.369% on sm_120) match. Occupancy table (regs 244/244/228/222, 67,584/101,376 B, 4/8 warps) matches.
• Frozen protocol elements: hardware + driver + CUDA toolchain + arch flags ✓ (L20 driver 580.126.20/CUDA 12.1; RTX 5060 Laptop 575.57.08/CUDA 12.9), base git sha = merge-base ✓, seeded inputs (mt19937 20261101) ✓ (synthetic, no checkpoint — declared), A-B-B-A interleaving/warmup/event timing ✓, byte-comparison fidelity criterion (no numeric tolerance) ✓, run budget + "nothing added after seeing results" ✓, stop condition ✓, acceptance criteria ✓. No tolerance appears in the identity claims. Body declares 512-row regression honestly and gates on it.
• Byte-identity evidence table (body "Fidelity" + README + results.json parity/sanitizer fields) present; sanitizer 0 errors at 4 lengths × 2 devices; prefix windows q_base 64/128/1024 vs full pass.

Findings

• P3 — readiness gap: zero executed CI runs at head (details above). The cuda-compile job has never seen this .cu; compile risk is author-attested only. Trigger: workflows awaiting maintainer approval (first-time contributor) + fork Actions disabled/empty. Consequence: cannot gate merge-ready on automated evidence yet; approving the workflow runs is the single highest-value next action.
• P3 — scope statement narrower than code: bidirectional full passes also take the wide tile. launch<true, true> (StartLux Instant cs1_attention_gated_bidirectional) hits the wide branch when q_base=0 and rows ≥ 640 on sm_89/sm_120, but the body scope bullet and both READMEs say "causal full pass and 128-aligned prefix windows". Structural bit-identity covers it, and the (GPU-gated) bidirectional reference test would exercise it at t=700/2048, but plan.json's measured/sanitizer arms do not mention the bidirectional mask. File: src/backends/cuda/qwen3_5/attention.cu, wide branch in launch<> (~line 316 at head). Fix: one sentence in body/READMEs, or exclude Bidirectional from the gate.
• P3 — process-wide static for capability decision. wide_tile_supported() caches on the device current at first call (attention.cu:292–299); on a heterogeneous multi-GPU host, the first-call device's capability applies to all later launches. Bounded consequence: worst case is the wide tile on an unmeasured arch (still structurally bit-identical; sm_90 would accept 101,376 B), or never selecting it; a mismatched cudaFuncSetAttribute on a second device fails loudly via the returned error. The same first-call-device pattern already exists (and is commented) for the 64-query path, so this is consistent-with-existing, not a regression.
• P3 — the cited "repository's guidance against per-PR microbenchmarks" is not findable on main. Tree-wide search finds no such written guidance; the nearest text (.agents/skills/self-review/SKILL.md ~line 140) actually asks to "Link raw repeated-run results and the protocol", and the closest precedent (benchmarks/gdn/artifacts/20261003/) committed harness programs and raw batches. The final trim (medians + delta ranges only) is more reviewable and the protocol/results are self-consistent, but the citation justifying the trim is imprecise; the owner should either bless the trim or ask for raw artifacts.

Notes (informational)

• Byte-identity/occupancy tests are GPU-gated #[ignore] tests, not CI-wired — pre-existing repo design (Tier 3 manual); on this PR it means identity evidence is author-run only until someone with a card re-runs cargo test ... -- --ignored.
• Body phrasing "cs1_attention_gated_shared, the batch entry points and every other architecture keep 64 queries" is loose: cs1_attention/cs1_attention_gated — the measured public entries — do widen when the gate passes; batch/decode calls stay on 64 only via the ≥640-row gate, and no exported symbol is named "batch".
• First commit carries 4,434 lines (raw jsonl ×~2.5k, logs, pr-description.md) that are all removed by head — fine for history; noting for completeness of the iterative-cleanup story.
• Occupancy math verified: 244 regs × 256 threads = 62,464 ≤ 65,536 and 101,376 B smem → exactly 1 block/SM → 8 warps, as recorded.
• Local (no GPU/cargo here): pytest tests/benchmarks -q 133 passed; mkdocs build --strict exit 0; worktree removed.
• No CI runs were observed anywhere for this PR (see Observed CI). Everything above labeled author-reported comes from the body/record, not from runs I executed.

Verdict rationale

DRAFT PR → verdict COMMENTED regardless. The kernel change is minimal and structurally sound (identity by unchanged per-row arithmetic, not by tolerance), the record is complete against the frozen-protocol checklist, the final tree is hygienic, ABI untouched. Nothing blocks entering review; the four P3 items are readiness/scope-documentation gaps to resolve before marking ready: approve the workflow runs so CI (esp. cuda-compile) actually executes, and reconcile the bidirectional scope sentence.

[psv666]

[P2] Key the new architecture decision by the current CUDA device.

wide_tile_supported():293–301 caches the first device's capability in a process-wide static bool. The backend exposes cs1_set_device, but after a first call on sm_89, a call on sm_90 still returns true and enters the wide path, contrary to the stated restriction to sm_89/sm_120. The reverse order permanently disables the optimization on a supported device. This also affects threads with different current devices.

I compiled the exact helper with CPU stubs for the two CUDA queries and changed the simulated current device from sm_89 to sm_90:

sm89 => 1
sm90 => 1

Please query the current device per dispatch or cache by device/context. The newly added wide_configured attribute setup at lines317–319 should follow the same device ownership; a process-wide successful setup does not configure a second device. A mocked alternating-device regression can cover the decision without heterogeneous GPUs. This is a host-dispatch finding, not an observed GPU numerical mismatch; a process that stays on one device does not trigger it.

I inspected the complete 5-file diff, unchanged key-tile/row arithmetic and prefix-alignment fallback. 26 backend CPU tests and strict MkDocs passed. No CUDA build, sanitizer or GPU parity was independently run. The author-reported GPU measurements remain separately scoped; I confirmed the older a3bb8c4 evidence commit includes the harness sources and raw records, so the current plan's “programs ... not committed” wording is stale.

Review scope and provenance

Head e9aeb07e2acc184bcc0f42dba5ed09121d99c850; PR API base 5fa5179d57f9267bb0d20162004d8ce26352aa9d; diff merge-base 5fa5179d57f9267bb0d20162004d8ce26352aa9d. Inspected scope: complete small diff, dispatch callers and historical evidence inventory. Complete changed-file inventory: 5 files, 57 changed authored-code lines / 232 total text lines (binary fixtures counted as files, not text lines).

Loaded trusted InferMatrixLab/system1-omni@5fa5179d57f9267bb0d20162004d8ce26352aa9d:.agents/skills/system1-omni-review/SKILL.md, SHA256 3a4b876162da14eee8db56af86664576aa5755b8f3e80b3072e944018647c3bc, plus architecture, repository contracts, CONTRIBUTING, self-review artifact requirements and CI/Docs workflows. Personally executed checks are distinguished above from author-reported GPU evidence. No GPU inference was run in this review; no approval or readiness transition is implied.

[sysuTobo]

Fixes for the two review rounds:

Round 1 (hsliuustc0106)

• Scope wording: the bidirectional full pass also takes the wide tile (shared launch<>, q_base = 0). READMEs and the PR body now say "full pass (causal or bidirectional) and 128-aligned causal prefix windows".
• Removed the claim of repository guidance against per-PR microbenchmarks; it is not in the tree. The record now points at the per-batch files in commit a3bb8c4.
• CI: still action_required on both workflows; I cannot approve runs as a first-time contributor. Local compiles: nvcc 12.9 -arch=sm_120 and 12.1 -arch=compute_89,code=sm_89.

Round 2 (psv666, P2)

• wide_tile_ready() replaces the process-wide statics: the capability check and the shared-memory opt-in are cached per thread and re-evaluated when cudaGetDevice reports a different device, so cs1_set_device and mixed-device hosts do not inherit another device's decision.
• CPU-stub alternation after the fix: sm_89 -> wide, sm_90 -> narrow, sm_120 -> wide, back to sm_90 -> narrow; the opt-in is issued once per device.
• No measurable dispatch cost: interleaved pre/post A/B through cs1_attention_gated at 1,024 and 4,096 rows gives -0.33% and +0.32%.
• plan.json no longer says the programs are uncommitted; it points at a3bb8c4.
• GPU run after the fix: attention kernel tests 7/7 pass.