Repository navigation
[FlyDSL] Raw-dialect cleanup (others): gdr_decode / chunk_gated_delta_h / pa_mqa_logits_fp4 - #4609
Merged
Merged
Conversation
vector.BroadcastOp/from_elements/FMAOp -> fx.Vector.filled/from_elements/fx.math.fma (keeps fused v_fma); arith.constant -> fx literals, drop arith import. rocdl.exp2 + gpu.shuffle kept raw (fx I/O). gfx950 byte-exact (16/16), perf within noise. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
llvm.inttoptr<3> -> fx.to_llvm_ptr; drop redundant _to_raw at scf.for range(init=) boundary. rocdl.mfma/ds_read kept raw. gfx950 byte-exact (15/15); lowered IR md5-identical (perf-neutral). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
create_buffer_resource+buffer_load/store -> make_buffer_tensor + copy atoms / indexed fx pointers; 1-writer scatter -> guarded fx.ptr_store; ArithValue -> fx operators. rocdl.mfma_scale/ds_bpermute + mlir_undef (poison pad) kept raw. gfx950 op tests pass; byte-exact (ISA+LLVM-IR identical); prefill ~11% faster. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Contributor
🏷️ CI GuideRuns automatically on every PR:
Extended tests (opt-in via labels):
|
- _i32_buffer: replace inttoptr(Int64(ptrtoint(get_iter))) retype round-trip with recast_iter; drop redundant _I32_MAX_RECORDS (make_buffer_tensor defaults to max_size=True = descriptor 0xFFFFFFFF). - prefill: collapse the duplicated cta_info buffer (cta_info_bt + cta_info_flat over the same ptr) into one width-4 buffer; read fields 4,5 via 2D indexing. Fold the per-CTA base with add_offset (element offset) instead of int byte math. - port copy_atom_call -> fx.copy for the vec4 KV load and the Q/QS/W loads. Validated byte-exact (cos=1.0) on decode + prefill op-tests, perf on par. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
_i32_buffer / _load_vec4_i32 / _pack_i32_pair_to_i64 / _pack_lo_i64x2_to_i32x8 were duplicated verbatim in the decode and prefill kernels. Move them to one pa_mqa_logits_fp4_common module and import from both (-88/+63 net, single source of truth). Drop the now-unused _llvm import from both kernels. Validated byte-exact (cos=1.0) on decode + prefill op-tests. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…kage Move fp8_mqa_logits, pa_mqa_logits_fp4, _prefill and _common into a dedicated kernels/mqa_logits/ package. Update the package __init__, the two op-test deep imports, and fp8_mqa_logits's sibling import (.tensor_shim -> ..tensor_shim). Validated: fp4 decode + prefill op-tests PASS (cos=1.0); fp8 module imports (kernel is gfx1250-only, skipped at runtime on gfx950). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…style) Replace the raw rocdl.mfma_scale_f32_16x16x128_f8f6f4 calls in the decode and prefill MQA-logits kernels with fx.make_mma_atom(cdna4.MFMA_Scale(...)) + fx.gemm(atom, c, a, b, c, scale_a=, scale_b=). Q/KV operands are now i32<4:1> register fragments (drops the manual v8i32 hand-packing), accumulators bridge the software-pipelined SSA carry via c_frag store/load, and e8m0 scale words are passed as plain i32 values. Prefill keeps its per-nt opsel_b via one atom per nt. Not ported to a tiled MMA on purpose: per-block e8m0 scales + prefill's per-nt opsel can't be expressed by a single make_tiled_mma (one atom / one opsel per tile); the only in-tree scaled-MMA precedent, mxmoe_gemm_v2, uses per-call fx.gemm for the same reason. The now-unused v8i32 pack helpers are removed from pa_mqa_logits_fp4_common. Validated byte-exact (cos=1.0) on decode (all shapes) + prefill (all shapes); decode perf on par (within run-to-run noise). ruff F401/F811/F821 clean. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Resolves the fp8_mqa_logits.py conflict: main's #4606 tiled-copy cleanup rewrote the file (dropping _to_raw) while this branch moved it into the kernels/mqa_logits/ subpackage. Kept main's body, re-pathed the shim import to '..tensor_shim'.
- black: the refactor left single blank lines before top-level defs in
pa_mqa_logits_fp4{,_prefill}.py, and the longer mqa_logits module path
pushed an import in the op-test past the line limit. Base was clean, so
these would have failed the format check.
- gdr_decode: fx.math.fma already returns an fx.Vector, so the
fx.Vector(...) re-wrap before .reduce() is a no-op. Dropped in 3 places.
The hoisted 'from .pa_mqa_logits_fp4_common import ...' sits directly against the third-party import block; isort wants a blank line between the third-party and local-folder groups. Both lines are PR-introduced, so reviewdog's diff_context filter failed the ruff job. Verified with the CI-pinned ruff 0.16.0 and black: both clean.
coderfeli
force-pushed
the
flydsl-cleanup-others
branch
from
August 8, 2026 08:51
e726adf to
9635490
Compare
waqahmed-amd-fi
added a commit
that referenced
this pull request
Aug 8, 2026
#4609 removed the `vector`/`arith` imports the per-channel branch still used, and the merge left `r_g_vec` outside the scalar `else`, so building a `gate_mode="kda"` kernel raised NameError. The benchmark's direct kernel call also needed #4573's q/k/v strides and read/write index split. Scalar path bit-identical; 47 tests pass on gfx950.
junhaha666
pushed a commit
that referenced
this pull request
Sep 8, 2026
…hem (#5289) * validate-kernel-pr: record non-SystemExit failures in the execution receipt `main()` caught SystemExit only, so a target dying on an uncaught AssertionError -- the standard shape of aiter's op_tests -- left `status` at 0 and the `finally` block wrote a receipt claiming `pytest_exitstatus: 0` before the exception propagated. The process exit code was right; the evidence on disk was not, and a red run was indistinguishable from a green one to anything reading the receipt. Reproduced with two script targets, one passing and one asserting: before ok exit=0 receipt=0 fail exit=1 receipt=0 after ok exit=0 receipt=0 fail exit=1 receipt=1 Re-raising keeps the traceback and the exit code intact; SystemExit(3) still records 3. * validate-kernel-pr: count untracked files as kernel changes `kernel_files_changed` came from `git diff --name-only HEAD`, which never lists untracked files, and the script has no `git add`/`--intent-to-add` anywhere. Files that `git apply` introduces as new therefore did not exist as far as the test-policy analyzer was concerned: a PR whose whole point is a new kernel read as touching no kernel code, and a tolerance widened alongside it was reported under the weaker finding. Reproduced on a scratch repo carrying one modified test and one added kernel: before changed=['op_tests/test_thing.py'] kernel_changed=False after changed=[..., 'csrc/kernels/new_kernel.cu'] kernel_changed=True `--exclude-standard` keeps .gitignore'd build output out; the index is untouched. * review-pr: derive the applicable rule set from the diff, and split the Tier table Step 3 was a 20-type checklist mapping onto 44 rules, applied by the model to itself. A model asked to read all 44 attends to none of them, and Step 1 already had the diff in hand and was already running Python over it, so the classification costs nothing to derive structurally. triage.py has two modes. `rules` emits only the families the diff triggers. `evidence` answers the question those families actually turn on: for a removed guard or a changed signature, how the symbol is handled ON HEAD -- fetched, rather than left to a prose reminder to go and grep. That reminder was already in Step 1 and was read and not acted on, producing a withdrawn `q_out is not None` finding on aiter#5143; `q_out` is std::optional on both sides, every call site is `has_value() ? data_ptr() : nullptr`, and the kernel guards `if(is_q && q_out != nullptr)`. The collector puts those three lines in front of the reviewer. Measured over 597 open aiter PRs (every open PR at the time, diffs fetched from the API): rules read median 12 of 44 (27%), max 29, none needed the full set coverage 597/597 matched at least one family precision most common family fires on 48%, none above half Derivation is conservative -- a family it cannot decide structurally is included, never dropped -- and it degrades rather than failing: an unparseable diff, an empty diff, or GitHub refusing to serve one over 20000 lines all fall back to the full 44-rule set, which is the pre-derivation behaviour. D9 is deliberately absent from the emitted sets. scan_index_width.py already decides that family structurally and its candidate list is the input to the finding; re-reading the prose adds nothing and costs attention the scanner-less families need. Step 4's Tier table listed `aiter/ops/*.py (any)` as Tier 1. By the table's own Q1 test that is defensible -- `__init__.py` does `from .ops.xxx import *`, so breaking any one of them breaks `import aiter` -- but there are 200+ files under aiter/ops/, and putting every single-kernel wrapper in the same tier as jit/core.py empties the tier of meaning: measured across the same 597 PRs it fires on 71% of them, and a mandatory full assessment demanded of 71% of PRs stops being performed at all. Tier 1 is now jit/core.py and __init__.py, whose failure mode is every op; a wrapper is Tier 3, whose blast radius is its own op, and the import-chain risk it does share is covered by the B6 export check attached to the ops-wrapper family. The full Step 4 assessment now fires on 6%. * review-pr: sweep the diff's added imports against the merge target `triage.py` grew a `symbols` mode that resolves every first-party import a diff adds, but nothing invoked it and the SKILL never mentioned it, so the code was dead. Step 1b now runs it against a fresh checkout of the branch the PR merges INTO -- not the PR base, and not whatever the reviewer has on disk. The root is the whole check. aiter#4994 added `from aiter.ops.flydsl.utils import is_flydsl_available`; #5116 (3b2a9ce6) had deleted that module from main 19 hours earlier. Resolved against #4994's own base the import is fine and the sweep returns clean; resolved against main-at-merge-time it reports. #4994 merged green (83571d9b) and was reverted 5.5 hours later (283e1d4b). A stale root reproduces exactly the silence that let it through. Two corrections that fall out of that: - The finding is a REBASE signal, not an accusation of invention. Renamed HALLUCINATED-SYMBOL to UNRESOLVED-IMPORT; a reviewer who reads "hallucinated" about an import that was valid when written will dismiss the whole class. - `from pkg import sub` binds a submodule, which need not appear in the package's `__init__.py`. Searching only the __init__ text called four real imports invented across the 600-PR corpus (aiter.jit/core, aiter.jit.utils/chip_info, flydsl.kernels/{buffer_ops,vector}). Report rate over that corpus drops from 14% to 10%, and every module-level hit that remains is a module absent from today's main. Verified: reports on #4994 against main, silent against the pre-deletion tree, silent on #5143. * review-pr: pin every step to the PR's own commits, not the reviewer's checkout Three ways the same PR produced different results, or no result, depending on who ran the skill and what they happened to have on disk. 1. `$BASE_SHA` and `$HEAD_SHA` were referenced in Step 1b and assigned nowhere. Under `set -u` that aborts the block, so the cross-file evidence collector -- the thing added specifically because prose telling a reviewer to go and grep was read and not acted on -- never ran for anyone, and neither did any step after it. Both are now assigned once in Step 1 from the API: the base branch tip already fetched into base_head.txt, and headRefOid from the PR metadata. 2. `git fetch origin ...` assumed the remote named `origin` is the repo under review. Anyone who has pushed a branch has a fork there instead, and `origin refs/pull/N/head` then resolves against the FORK's pull refs: a different PR's head, or nothing. Every fetch now names $REPO_URL explicitly. 3. The evidence collector prefixed changed paths with $PROJECT_ROOT, reading the reviewer's WORKING TREE -- whatever branch is checked out, plus uncommitted edits -- rather than the PR's head. It now materialises head blobs under $WORK and names the ones it could not get. Measured on the 600-PR corpus, resolving against a checkout one day stale rather than the base tip changes the symbol sweep's answer on 33 PRs, 31 of them missed reports -- the aiter#4994 failure mode, arrived at silently. Verified end to end against live PRs (#5240/#5242/#5244 clean, #4166/#4441 report their flydsl.utils import) from two checkouts: current main, and one 130 commits stale with a fork origin, a dirty tree, and the deleted module still present. Byte-identical output from both. * review-pr: make the rule pass produce an artifact the skill can check Step 5 asked the model to work 44 rules -- now a derived 12 -- and marked them off to itself. Nothing downstream could tell a run that adjudicated every rule from one that read the list and went straight to the card, and the skill already has its own evidence that the second is what happens under load: on a 14-PR controlled run the revised D9 text caught 0 of 3 known overflow defects and the scanner it names was never invoked once. That is why the D9 scan was moved into Step 1. This applies the same move to the rule pass itself. Step 5 now writes one line per derived rule to $WORK/verdicts.txt: <RULE-ID> FIRE|CLEAR|N/A — <reason naming file:line, symbol, or condition> and Step 8 will not write a card until `triage.py ledger` is green. The gate reports UNADJUDICATED for a derived rule with no line and NO-EVIDENCE for a verdict with no reason, and fails closed on an empty rules.txt -- the state of a run whose Step 1b never happened, which must not be the cheapest path to a verdict. It checks that each rule was answered, not that the answer is right. Answering wrongly is a review error a human can see in the card; not answering is invisible, and invisible is the one that scales. Verified against the four shapes of skipping: no Step 1b (exit 1, LEDGER UNUSABLE), card without a rule pass (exit 1, 4 UNADJUDICATED), two rules cherry-picked (exit 1), every rule stamped "ok" (exit 1, 4 NO-EVIDENCE). A real adjudication passes 4/4. * review-pr: give Triton kernels their own rule families Triton is 195 of the 600 open aiter PRs in the replay corpus, and the derivation said almost nothing about them: `ops-wrapper` fired on 88% and `modified-kernel` on 79%, so a Triton PR arrived labelled rather than triaged. The generic families were tuned on the whole corpus, and this is the subpopulation where that shows. Four families, mined from what the 195 diffs actually contain, each firing on at most 52% of Triton PRs and 17% of all of them: triton-mask-bounds T1 T2 27% unmasked tl.load/tl.store; mask= without other= triton-launch-cfg T3 T4 52% num_warps/num_stages/num_ctas; waves_per_eu, matrix_instr_nonkdim, kpack across archs triton-accum-prec T5 25% tl.dot, .to(tl.float32), allow_tf32 triton-grid-map T6 27% host grid vs the kernel's program_id mapping `tl.load`/`tl.store` arguments are extracted paren-balanced rather than by regex: `[^)]*` stops at the first close paren, so `tl.load(p, mask=(a<b), other=0.0)` reads as having no `other=`. Twenty of the 195 contain that shape; the balanced scan drops the mask family from 55 PRs to 52, all three of them false. T1/T5/T6 are 🔴-eligible and carry their own FP self-checks -- an unmasked load is fine when the axis is a constexpr multiple of the block, and T3 is a perf finding that needs a base-vs-head number or it is `[inferred]` per Step 8. Full-corpus replay with the six new rules in: median 13/49 (27%, unchanged -- the additions did not dilute), coverage 598/598, no family over half. * review-pr: anchor a FIRE verdict to a file the PR actually changes The ledger checked that every derived rule had a reason, which closes "did not answer" but not "wrote something plausible". A reason is free text and mostly unjudgeable, but one part of it points outside the model: the path it cites. `ledger` now takes the diff and rejects a FIRE whose reason names a file this PR does not touch -- a defect claimed against untouched code is either carried over from another review or invented. Only FIRE is checked, deliberately. A CLEAR legitimately cites files outside the diff: "E5 CLEAR: aiter/__init__.py unchanged" is a correct reason precisely because that file is absent, and flagging it would teach the reviewer to stop citing anything -- worse than not checking. Reasons that name no file are not flagged. This narrows the gap between "answered" and "answered honestly"; it does not close it, and nothing here judges whether a CLEAR is correct. Verified: a fabricated FIRE citing aiter/fused_moe.py on a PR that changes only aiter/mla.py and op_tests/test_mla_persistent.py exits 1 with UNTOUCHED-CITATION, while the same ledger with the FIRE moved onto aiter/mla.py:120 passes 4/4, and the two CLEARs citing untouched files pass in both. * review-pr: take the rule bodies and Step 1 out of the entry file SKILL.md went 487 -> 632 -> 1372 lines between 2026-07-09 and 2026-09-03, and almost none of that was instructions. At 1210 lines, 825 were inside code fences and 299 were prose; Step 1 alone was 828 lines, 95% of it bash. It was a ~300-line skill wrapped around a 785-line shell script, and a reader had to scroll past the script to reach the part written for them. Two moves, same principle: what is conditional gets derived, what is executable gets called, and only what every review needs stays resident. - Rule bodies move verbatim to rules.md, byte-identical to the span they came from. `triage.py expand` writes $WORK/rules_expanded.txt: the full text of exactly the derived rules, under their headings, nothing else. Cutting which rules the reviewer is TOLD to check from 44 to a median of 12 while still shipping all 44 texts was half a fix -- 383 lines of bodies against a median need of 103. Corpus replay: 0/600 expansion failures, median 117 lines against a fixed 410 (29%). - Step 1 and 1b move verbatim to fetch.sh, which SKILL.md calls in three lines and which prints its scratch dir for the steps that follow. SKILL.md: 1210 -> 440 lines, below the 487 it started at, with everything added since still in force. Two shapes the expander had to learn, both found by running it over 600 real PRs rather than by reading the file: Housekeeping rules are table ROWS, not `**HKn — **` blocks (parsing only blocks reported HK4/HK6/HK9 as undocumented on 277 PRs -- they are documented, the parser was wrong), and STEP4 names Step 4 rather than a rule body. A rule whose body genuinely cannot be found is a hard failure naming the ids, because emitting 11 of 12 bodies reads exactly like emitting 12. Verified equivalent, not assumed: fetch.sh and the previous inline blocks were run against aiter#5222 and #4166 from the same checkout. stdout matches line for line and pr.diff, rules.txt, rules_expanded.txt, symbols.txt, evidence.txt, validation_requirement.json and base_head.txt are byte-identical. The move broke four tests in validate-kernel-pr that assert on strings in review-pr's SKILL.md. The contracts are right -- perf-harness detection must agree across both skills, the scanner must be required, the identity gate must be the one shipped -- they were reading half the skill. They now read SKILL.md and fetch.sh together, and the gate block is selected by containing `expected_verdict` rather than by being the second heredoc, which is not a property anyone maintains. Mutation-testing those repairs found two holes. `assertIn("fetch.sh", step1_block)` was satisfied by the COMMENT above the call, so breaking the call site left it green; it now requires a non-comment line. And the identity gate had no test at all: every case fed it `headRefOid` copied from the report under test, so replacing `actual_head != expected_head` with `if False` -- a gate accepting a report from any checkout -- kept all 88 tests green. test_review_gate_rejects_a_report_for_another_head feeds a mismatched head and is confirmed to fail against that mutant. rules.md and fetch.sh are force-added: .gitignore line 21 ignores .claude/ wholesale, and everything already tracked under it was force-added too. Without -f these two files stayed untracked while SKILL.md was committed calling them. 89 validator tests, 241s. * review-pr: test what the deriver produces, and budget the entry file Mutation testing over triage.py: 170 predicate flips, 20 caught. `and` to `or`, a dropped `not`, `p.count("/") == 1` to `!=` -- 150 of 170 mutants left the suite green. Those predicates decide which rules a reviewer is shown, so any edit could silently change every review with nothing failing. The 600-PR corpus existed, but as something run by hand, not as a test. corpus.tgz pins 59 real PRs: 19 chosen so every one of the 31 families fires at least once, 40 more drawn with a fixed seed. expected.json holds their exact `rules` and `symbols` output. When a deriver change is intended, re-bless with bless.py and commit the diff -- that diff names which real PRs would now be triaged differently, and it is the review. Two invariants beyond byte-equality: every corpus PR derives at least one rule (an empty result silently clears a PR), and median expansion stays under 45% of rules.md (the derivation can stop discriminating while every individual output still looks reasonable). That took the score to 38%, and the survivors clustered. A third sat in sweep_symbols/symbol_defined/resolve_module, because only 6 of the 59 corpus PRs emit any UNRESOLVED-IMPORT line -- the paths that decide NOT to report were barely reached, and those are the ones that fail silently. Nine cases cover them directly: namespace package, submodule binding, star-import and __all__ re-export, module added by the PR itself, symbol added by the PR itself, deleted-import lines, third-party and relative imports, and the two genuine misses. Every one had been hand-verified when written and then not encoded. Eight more cover check_citations, where seven mutants survived: nested paths, no line number, a FIRE naming no file at all (demanding one would push the reviewer to invent a citation), CLEAR and N/A citing untouched files, and the two-argument call with no diff to judge against. The structural half holds the shape the entry file was losing: SKILL.md budgeted at 500 lines, no code fence over 30 lines, every rule id the deriver can emit has a body in rules.md, every family expands to something, and fetch.sh's `required ... missing` guards must each be followed by a non-zero exit -- mutation turned all six of those into `exit 0`, so a missing scanner or schema would have let the run continue and report nothing, and nothing failed. That last one is a lint over the script's text, not an execution test; running fetch.sh needs a GitHub round trip. Both structural limits were confirmed to bite by pushing the file past them. Without the budget this comes back: nothing stopped the file growing the first time, so it grew. Raising the number is now a commit that has to say why. 35 tests, 5.4s, stdlib only, no GPU and no network. * review-pr: close the gaps found by reviewing ten real PRs Reviewing a random ten of the 619 open aiter PRs, one at a time, the way the skill is actually used. Two PRs carry defects that would land: #5220 inserts two ints into the middle of a kargs struct whose layout is pinned by `static_assert(sizeof == 112)` plus one `offsetof` per field, in the same translation unit it edits, so all three assertions fail; #4146 imports `aiter.ops.triton.utils.core`, which #5061 renamed to `utils.config_utils`, so it is a ModuleNotFoundError on merge. Seven open PRs still import that module and 41 still import the deleted `aiter.ops.flydsl.utils`. Five gaps in the skill, each found by using it rather than by reading it: - **No rule covered struct layout at all.** D11 fires when a struct in C-like source gains or loses a field and the diff touches no `offsetof`/`static_assert(sizeof` line -- derivable purely from the diff, 3% of the corpus, 40 such assertions across 37 files in csrc/. It does not fire when the PR updates the table in the same diff, which is the correct shape. The struct declaration usually lives in the hunk header's scope context rather than the hunk body, which is why detection reads `@@ ... @@ struct X`. - **The sweep skipped relative imports entirely.** `from .chip_info import ...` inside aiter/jit/utils/asm_guard.py resolves to aiter.jit.utils.chip_info. Skipping them left the case where an unresolvable import is worst: a Tier-1 PR wiring a new module into aiter/__init__.py, where it breaks `import aiter` outright. #5107's chain is sound, but only because it was checked by hand. Corpus report rate 9% -> 13%, all of it the same deleted modules reached through a second syntax. The `mod.startswith(".")` skip is now dead and removed. - **The sweep read non-Python files.** docs/tutorials/add_new_op.rst walks a reader through writing an op, imports and all, and two PRs were reported for a module named in a tutorial. - **HK9 fired on benchmark scripts.** PROF_WARMUP/PROF_ITERS in op_tests/flydsl_tests/profile_flydsl_bwd.py are not permanent runtime knobs; HK9 exists because aiter reverted one that was. The family now reads runtime files only. - **The evidence collector ignored the most common shape of the finding it serves.** The invariant-removed family fires on a deleted `.contiguous()`, but the collector only extracted symbols from assert/*_CHECK lines, so a removed contiguity normalisation produced no evidence and the reviewer was back to grepping -- which is what the collector exists to stop. #5235 relocates two of them; the answer is "it moved", and now the collector shows it. Also generated: Step 3's type -> rule table moves to MAPPING.md, emitted by `triage.py mapping`. The hand-written copy had drifted -- it claimed D9 was derived when D9 is scanner-backed and deliberately is not, and omitted 21 rules that are, including every Triton rule. A mapping the reviewer is told not to apply by hand should not be maintained by hand. Step 6's first structural check asked the model to list every new symbol and grep it; `triage.py symbols` had already done that, so it now points at $WORK/symbols.txt and scopes the manual work to what the static sweep cannot judge. The other five checks write one line each into $WORK/ai_diagnostic.txt, and Step 8 runs answers, diagnostic and ledger before a card may be written. 600 PRs: median 15/50 rules (30%), 600/600 covered, no family over half, symbols 13%. 99 structural tests, 89 validator tests. Mutation score on triage.py 93% of 137 code mutants; the nine survivors are equivalent mutants, listed and not papered over with tests that cannot kill them. * review-pr: put the redundancy and test-quality evidence on screen Asked why a ten-PR pass reported no redundant code, no AI slop and no weak op_tests, the answer was that none of those has a detector. Step 6 asks for all three in prose -- "identify mirrored code ... compare it field by field", "is the reference impl structurally a twin", "is atol/rtol loosened with no justification" -- and prose asking for a search is the shape that does not happen, which is why D9's scan moved into Step 1 and why the rule pass grew a ledger. Two collectors, both evidence rather than verdicts, because measuring first showed a verdict would be wrong more often than right: - `twins` names the file each new file was copied from, with the overlap ratio. 3% of the 600-PR corpus adds a file at least 60% identical to one already in the tree: fused_gemm_a16w16_copy_x.py against _quant_x.py at 65%, test_opus_gmem_gfx1100.cu against _gfx1201.cu at 75%. The pair is not the finding -- an arch-specific variant is the normal shape -- the asymmetry between them is, and that needs a human diff. - `testquality` prints, per added test file, the number of assertion primitives, the tolerances and the shapes. It does NOT conclude "this test asserts nothing", because that is not decidable from a diff: assertions routinely live in a shared helper (`run_fp8(..., verify=True)`), and the first two corpus cases inspected under a firing rule were both false. The strict form -- a whole new test file with no assertion primitive anywhere -- fires on 2 of 600, one of which is a benchmark. A zero count is printed with the reason it is not yet a finding. Measured and rejected as rules, recorded here so they are not re-proposed: intra-diff six-line duplication (30% of PRs, no discrimination), test bodies with no assert (over half false), and a reference implementation calling the op under test (1 of 600). 112 structural tests, 89 validator tests. * review-pr: HK12, a test file pytest collects nothing from Batch two. aiter#4821 ships op_tests/flydsl_tests/test_dispatch_tdm.py and compile_dispatch_tdm.py: both named like tests, both written as `main()` scripts with `_check()` helpers, neither containing a single `def test_`. pytest imports them and collects zero tests. HK6 asks a new op to ship `op_tests/test_*.py` and this satisfies it by name while contributing nothing to CI. The rule is deliberately narrow. aiter's tree already holds such files, so the style is tolerated and the file is not itself a defect; HK12 fires only when the PR ships no collectable test at all -- the case where HK6 reads as satisfied and the new code is in fact untested. 3.8% of the 600-PR corpus, against 4.3% that add such a file at all. The evidence collectors added for batch one earned their place immediately on batch two, in both directions: - test_quality flagged `atol, rtol = 0.125, 0.125` in aiter#3940, and the four lines above it in the diff derive that number from e4m3's three mantissa bits. Not a finding, and it took one look rather than a grep -- which is what putting the number on screen is for. - test_quality reported `test functions added : 0` for #4821's two files, which is what led to HK12. Median rules 15/51 (29%), HK12 3.8%, 117 structural tests. * review-pr: check whether a new test will ever run, not whether it looks like one Challenged on the HK12 finding from batch two -- "op test CI should pick all of those up, did you judge it wrong" -- and the answer is yes, the mechanism was wrong and the real one is worse. .github/scripts/split_tests.sh shard-scans `op_tests` with `-maxdepth 1` and `op_tests/triton_tests` recursively. Nothing else is scanned by anything: op_tests/*.py 127 files aiter shard, runs by default op_tests/triton_tests/** 107 triton shard, runs by default op_tests/opus/ 1 opus-test.yaml op_tests/multigpu_tests/ 29 only with the `multigpu` label, skipped by default op_tests/flydsl_tests/ 5 no workflow scans it at all So aiter#4821's two files would not run even if every function in them were a `def test_`. The absence of `def test_` was a real observation about the wrong thing. `triage.py citest` answers the question that matters -- will a CI job run this file -- and answers it from .github/ rather than from a table here, because a copy of the shard script would drift the way Step 3's rule table did. 9% of the 600-PR corpus adds a test that never runs: 48 into unscanned directories, 6 label-gated. One trap on the way: a path MENTIONED in a workflow is not a path that gets RUN. pr-title-tags.yaml names op_tests/flydsl_tests/ to choose a label, and counting that as coverage reported #4821 as covered -- the exact case the collector exists to catch. Only lines that execute something count now, and a test pins that. HK12 is now about location and is evidence-backed, listed in UNREACHABLE_BY_DESIGN alongside D9 since the diff alone cannot answer it. The original observation survives as HK12b at 📝, explicitly secondary: check where the file lives before asking what is in it, because a file in an unscanned directory does not run whatever it contains. The unreachable-rule invariant caught this during the edit -- HK12 documented with no family emitting it -- which is what it was added for. Lettered HK ids (HK12b) needed the id pattern widened; the expander was silently dropping them. 124 structural tests, 89 validator tests. * review-pr: list the perf claims that name no baseline P1 asks for a number with its units and its comparison, and P6 asks for base-vs-head. Neither had anything reading the description, so a claim with no other side to it reached the reviewer as prose. `triage.py perfclaims` reads the PR body, lists every numeric claim, and marks the ones that say nothing about what they are measured against. Four things it deliberately does not flag, each found by running it rather than by reasoning about it: - Markdown table rows whose header says `before | after | speedup`. aiter#4443 ships six such rows and judging each row alone called all six unbaselined. - Signed deltas. `+8.64% end-to-end` compares against not having the change, which is a stated baseline in ordinary English; interrogating it is the pedantry that makes a check ignorable. - Shares. `33.07% of GPU time`, `81.8% prefill by GPU time` say where the time goes. - GPU counts. `8x MI355X` is eight cards. Substituting the model name away first left a bare `8x` reading as a speedup. Scoped to the description on purpose: reading claims out of kernel comments over the 600-PR corpus was 61% noise -- `4x DS_READ`, `<4 x i32>`, `num_tokens x 384 x 7168`, `5% of elements`. The claim lives in the description, which is where P1 and P3 ask for it. Measured and rejected in the same pass, recorded so they are not re-proposed: - Comment density. Median 7% of added code lines across 409 PRs, P90 21%. The five PRs over 35% were all read: aiter#4061 repeats one explanatory line across 140 generated CK files, #3557 is a section banner for the RDNA4 WMMA path, #2814 explains non-temporal loads. High density tracks codegen and documented kernels, not slop. - "Comment restates the next line", 12% of PRs. Inspected: it is mostly matching two adjacent comment lines, and the one real comment-then-code case was a useful comment above an assert. 132 structural tests, 89 validator tests. * review-pr: D11 asks the tree whether a layout is pinned, not the diff Batch five caught the rule added in batch one being noisy. aiter#5223 changes six headers and D11 fired, because the family was "a struct field moved and no assertion line changed" -- which is true of any struct field churn anywhere. None of those headers has an assertion within reach. A layout change is a defect only when something asserts the layout, and that is a fact about the tree. `triage.py structabi` extracts the struct names whose fields moved, then looks for `offsetof(<name>` or `static_assert(sizeof(<name>))` in the repository, and reports only the ones it finds. Corpus report rate 3% -> 0.3%: two PRs, both touching `pa_sparse_prefill_kargs`, which is the struct that actually carries the table. The true positive from batch one, aiter#5220, survives the tightening. D11 joins D9 and HK12 in UNREACHABLE_BY_DESIGN, all three for the same reason: the diff alone cannot answer the question, so a Step 1b collector answers it and the rule reads the artifact. Also from batch three: BASELINE was English-only, and aiter's descriptions are partly Chinese. aiter#5043's table is headed `场景 | 输入数据 | 时间 | TFLOPS` and correctly flagged -- it reports absolute times for a new ASM kernel whose own description says the shape previously fell back to CK/Triton, with no measurement of what it replaces -- but a table headed `优化前 | 优化后 | 提速` would have been flagged too. Both directions pinned. 136 structural tests, 89 validator tests. * review-pr: reduce a comment-only diff to the lines that are not comments Batch nine. aiter#4062, "docs(python): condense verbose comments", is 252 files and 7963 changed lines. A reviewer either skims it or spends hours on it, and both of those miss whatever code is hiding in it. The answer turns out to be: nothing. It is worth being able to say that in one line. `triage.py commentonly` fires only when a diff is at least 90% comment churn and at least 60 lines -- 2 of the 600-PR corpus -- and prints the changed lines that are not prose. Three passes were needed before it told the truth about #4062: - Trailing comments are stripped before comparing. `REVISION = 25 # rev24, drop the token-major path` becoming `REVISION = 25 # g layout fixed head-major` changes no code, and #4062 read as four code lines until that was handled. - `/* ... */` continuation lines with no leading `*` are prose. aiter#4061's rewritten quant_utils.cuh header is exactly that, and every prose line in it was listed as something to review. - Those lines are marked as prose rather than dropped, because the floor and the ratio are about the whole diff; dropping them made a 6000-line comment rewrite look like a two-line diff that fell under the floor. Both PRs now reduce to zero code lines, which is the truthful answer. Same batch, a smaller correction: `invariant-removed` matched `assert` / `.contiguous()` in deleted COMMENT lines, so #4062 fired D4 on prose about asserts. Measured across the corpus this is 3 of 71 fires, so it was noise rather than a hole -- fixed because it is free, not because it was urgent. 144 structural tests, 89 validator tests. * review-pr: a fourth gate that asks whether each finding is nailed down answers, diagnostic and ledger all check that the work happened. None of them looks at the card, which is the one artifact a reader actually gets -- so a review could adjudicate 27 rules honestly and then write a finding that came from none of them. And the red threshold, "before firing any red, write down the concrete input that triggers it", was prose with nothing reading it. `triage.py card` checks three things per finding and only three: - UNTOUCHED-FINDING: it cites a file this PR does not change. - UNBACKED-FINDING: it carries a rule id no verdict marked FIRE, or names files that appear in no verdict, diagnostic or blind-spot line. - UNPROVEN-RED: a red naming neither a value nor two code identifiers. Whether a finding is correct is not checked and cannot be; that is what a human reads the card for. The concreteness test needed a second pass. Requiring a digit rejected a real finding -- "row=(tail_blk*num_kv_heads+kv_head_idx)*BLOCK_M against nrows=tail*num_kv_heads*BLOCK_M" is as concrete as a finding gets and contains none. Two distinct code identifiers is the line between that and "the reduce kernel looks racy". Adding the gate pushed SKILL.md to 512 lines and the budget test failed, which is what it is for. The budget was not raised: Step 1b had grown a paragraph per artifact, nine of them, each restating what the artifact prints about itself. They condense to a table -- file, what it answers, the trap it exists for -- and the entry file lands at 477, twenty under where it started this session and thirty under the budget. 152 structural tests, 89 validator tests. * review-pr: a relative import inside __init__.py resolves to that package Batch eleven. aiter#4515 was reported twice for the same module under two names, one of which does not exist anywhere: `aiter.ops.kernels.pa_mqa_logits_fp4`. The resolver stripped `.__init__` from `aiter/ops/flydsl/__init__.py` to get the package name and then dropped a component as well, which is right for `aiter/ops/flydsl/x.py` and wrong for the package's own `__init__.py` -- the containing package IS aiter.ops.flydsl there. Corpus report rate 13% -> 10%, 144 lines -> 104: fifteen PRs were being reported for imports that resolve fine one level up. The remaining finding on #4515 survives and is real: `pa_mqa_logits_fp4.py` was deleted by #4609's raw-dialect cleanup and the kernels live under `mqa_logits/` now, so the PR modifies a file that no longer exists on main -- a modify/delete conflict on merge, and imports pointing at nothing. That is the fifteenth distinct module broken by a reorganisation in the open queue. Same batch, worth recording rather than fixing: aiter#4848 fixes MoE expert addressing past 4 GB -- a 32-bit overflow class -- and ships op_tests/flydsl_tests/test_flydsl_moe_4gib_addressing.py, which lands in the directory no CI job scans. The test for a subtle addressing bug will never run. ci_coverage caught it, which is what it was added for. 154 structural tests, 89 validator tests. * review-pr: pull the subject out of a removed guard whatever shape it takes Batch twelve. aiter#4255 deletes `assert gfx_version in ("gfx942", "gfx950", "gfx1250")` and the invariant-removed family fired, correctly. The answer is one line away and was not on screen: the assert is REPLACED, by `assert _is_gluon_pa_mqa_logits_supported( gfx_version)`, which is a stronger guard than the allowlist it replaces. The evidence collector said nothing, so clearing it meant grepping by hand -- which is the thing the collector exists to stop. It extracted symbols from exactly two shapes, `X is not None` and `X->`. Across the corpus, 57 PRs delete a guard and 51 of them -- 89% -- produced no evidence at all: assert causal, "Only causal attention is supported" -> causal assert WQ.dtype == dtypes.fp8, "only fp8" -> WQ assert num_head_qo % 16 == 0 -> num_head_qo assert gfx_version in ("gfx942", "gfx950") -> gfx_version TORCH_CHECK(out.is_contiguous(), "out must be ...") -> out The subject of a guard is the first identifier after the keyword, skipping the words that are never a subject (`not`, `is`, `None`, `len`, `isinstance`, dtype namespaces). 89% -> 4%, averaging 3.6 symbols per PR. This is the third time the same shape of gap has appeared: a family fires on a pattern wider than the collector serving it understands. First `.contiguous()` removals, then comment lines, now every assert that is not `X is not None`. The family and its collector are written in different places and drift apart; both are now covered by cases that name the shapes explicitly. Rest of batch twelve was quiet: no unresolved imports, no twins, no unrun tests, no pinned-layout changes, no unbaselined claims. 156 structural tests, 89 validator tests. * review-pr: the card gate now runs in both directions, and fails closed Asked whether the "is this nailed down" step was actually nailed into the skill, the answer was no. Probing my own gate found three ways past it: - A missing $WORK/card.md read as an empty card, which read as no findings, which passed. Not writing the card was the cheapest way past the check. Now a hard failure. - A card with no findings passed while the ledger held nine FIRE verdicts. I had checked that findings trace back to work and not that work reaches the card, which is the inverse and the one that matters more: a review can adjudicate nine rules FIRE, report none of them, and look clean. - A card carrying only a heading passed the same way. FIRE now means it goes in the card. The 5-finding cap is a real reason to leave one out, so the escape is `-- not reported: <reason>` appended to the verdict line -- written down rather than silent. Verified against the real aiter#5157 ledger: nine FIRE, five on the card, two withdrawn in writing, and the gate holds out for the remaining two until they are accounted for. One test had to change rather than the product: a note citing aiter/mla.py with nothing in any verdict, diagnostic or blind-spot line mentioning it is a finding from nowhere, and Step 7.5 is where free-form observations get written down. Both directions are now covered -- a note grounded in the blind-spot answer passes, one from nowhere does not. 161 structural tests, 89 validator tests. * review-pr: a benchmark with no assertions is a benchmark Batch thirteen. One evidence line across ten PRs, and it was noise: aiter#4016 ships op_tests/triton_tests/bench_gdn_chunk_prepare.py beside a real test, and test_quality told the reviewer its zero assertions "may be in a helper, or there may be none". A benchmark measures; it asserts nothing by design. 68 of 682 rows over the corpus are that shape. bench_*/profile_* files and anything under op_benchmarks/ are now labelled, and their zero count reads "expected for a benchmark". The same count on a real test still asks about the helper, because there it is a real question. Batch thirteen otherwise clean: no unresolved imports, no twins, no unrun tests, no pinned-layout changes, no unbaselined claims, no comment-dominated diffs. While adding the tests, a behaviour worth pinning turned up: a test file with no tests, no assertions, no tolerances and no shapes is not listed at all. There is nothing to put in front of the reviewer, and that is right, but nothing said so. 165 structural tests, 89 validator tests. * review-pr: T8, a launch knob hardcoded past the config system aiter#5137 was reverted on one review comment: "Please change the config file jsons instead of hardcoding it in the code waves_per_eu=2". T4 already covers a knob forced across archs; this is the single-arch case and a different failure. The config system exists so a tuning sweep and the code disagree loudly, and a literal in kernel source is invisible to the sweep -- the next sweep either overwrites it or fights it, and neither shows up in a diff. #5137 carried a six-line justification for the value; the justification was not the problem, the location was. Fires on waves_per_eu / matrix_instr_nonkdim / kpack / num_stages / num_warps / num_ctas set to a literal in a .py outside configs/ and op_tests/. 11% of the 600-PR corpus. Exempt: a knob read from a config (`num_warps=config["num_warps"]`, `**cfg`), a literal in a config file, a test pinning a knob deliberately, and a commented-out line. Recorded with it: the pipeline is lenient. Measured across all 600 PRs, 77% produce no evidence line at all from symbols, twins, citest, structabi, commentonly or the test-quality flags. Some of that is correct -- most PRs are config bumps and tuning rows -- but 77% is not a number to be comfortable with, and the next passes should be spent raising recall rather than trimming noise. 171 structural tests, 89 validator tests. * review-pr: run the fail-closed guards instead of reading them * review-pr: Step 4 writes down which backbone files it cleared, and why * review-pr: stop tracking __pycache__ * review-pr: ask about a parameter that left, and a kernel that shipped untested Two of eleven candidates measured against the 600-PR corpus survived being read. api-signature required a `def`/`void`/`template` line on BOTH sides of the diff -- a signature rewritten in place. Deleting one line of a multi-line signature never touches the `def` line, so B6/E1/E5 were never asked about it. Reported as 9.3% of the corpus; 9.3% is what the shape scores before it is read. A hunk header names the ENCLOSING def, so `transpose_out=True,` passed to a call inside `def asm_moe(` counted as a dropped parameter of asm_moe, and the same line pattern over C++ reads `std::optional<Tensor>& fp8_out,` as a parameter named `std`. Walking the old side back to an unbalanced `(`, and staying out of C++, leaves 3.5% (21/600) at roughly 8 in 10 on a sample of ten. It moves the family from 17.5% to 20.3% of the corpus, and four PRs in the pinned 59 now derive it -- 1831, 2577, 2592, 4997, each a real kernel parameter. That is the whole blessed diff; symbols and MAPPING are untouched. HK6 already fired on new-kernel but was prose, so the reader established the absence of a test by hand. kernel_tests.txt does it: 6.5% (39/600) add a kernel with no test pytest would collect. Collectibility is decided by filename, which is the rule pytest uses -- an op_tests/-only whitelist called aiter#3991 untested while it added five test files under aiter/aot/flydsl/tests/, and excluding any path containing "bench" called aiter#2889 untested on the strength of test_rmsnorm_bench_against_aiter.py, which holds two collectible tests. An evidence line that states a falsehood is worse than one that says nothing, so the corpus is checked for that specific lie and reports none. Rejected, with the numbers, so it is not proposed again: added `except` blocks whose body is only pass/continue/return. 10.7% of the corpus, and reported at 4/4 true. 38 of the 64 PRs match nothing but a capability probe -- `except ImportError: return False` is correct code -- and six sampled from the narrowed pass/continue set were all deliberate: a barrier timeout whose comment explains that hanging is the alternative, `int(os.environ["MAX_JOBS"])` falling through to a default, a ROCm-discovery fallback chain, a best-effort rmtree, a tuning loop skipping unsupported configs. That is the shape of the weak no-assertion check this project already rejected: high rate, plausible story, false on reading. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * review-pr: why A1 still has no evidence behind it A1 asks whether a sibling kernel has the same bug and has never had a forensic answer, so the candidate for one was measured: deleted lines still present verbatim in a file this PR does not touch, three or more times. 13.2% of 600 open PRs, 10.0% once the line is required to carry an operation instead of being a declaration, a bare `for`, or a docstring. Eight read: about half are what A1 means -- variant MoE ops sharing `A_scale.stride(0) if A_scale is not None and A_scale.ndim == 2 else 0,` across moe_op_silu_fused.py and moe_op_gelu.py, where fixing one and not the other is the defect. The rest are boilerplate two files share for no reason: `for(int i = 0; i < 4; i++)`, an OptionalHIPGuardMasqueradingAsCUDA line, `HEAD_DIMENSION_OPTIONS = [128]`. Narrowing does not fix it. It silenced two of the four known-noise hits, left the stream guard and the config constant firing, and cost one of the three real ones. That is the shape of the problem rather than a tuning depth: a sibling that needs the same fix and a sibling that shares boilerplate emit the same evidence, because it is the same line. Verbatim identity is the wrong signal. A1's own example -- PR#3841, a decode kernel fixed while `_prefill_opt` beside it was not -- is same-FILE. Sibling FUNCTIONS within the changed file is the direction that has not been tried, and it needs its own numbers before it goes in. Recorded in the rule body so the next pass starts there instead of re-measuring this. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * review-pr: hand A1 the variant it is asking about A1 asks whether the sibling kernel carries the same bug and has never had a forensic answer. Its own example is aiter#3841 -- a decode kernel fixed while `_prefill_opt` beside it was not -- which is two functions in ONE file, and twins compares whole files. So the last attempt looked across files for a deleted line still present verbatim somewhere untouched: 13.2%, half of eight read were boilerplate two files share for no reason, and narrowing cost a real hit for every two noise hits removed. That is recorded in the rule body; this is the version scoped the way A1 means. Same file, and the two function names must share an eight-character stem. `_moe_gemm_a8w4_decode` against `_moe_gemm_a8w4_prefill`; `_stage1_ragged_k` against `_ragged_k`, still carrying `mask = context_idx + tl.arange(0, ChunkK) <= context_length`, which is the bounds check A1 names; `mla_decode_fwd` against `mla_v40_decode_fwd` over `and q.dtype == dtypes.fp8`. 18.0% of 600 open PRs. Without the stem, any two functions sharing a file qualify and it is 23.3%, mostly helpers that merely sit together. The stem is deliberately not a list of variant suffixes. Adding one -- _opt, _v2, _prefill, _decode, _fwd -- gives a tidier 10.8%, and 43 of the 43 it drops include kernel_unified_attention_2d against _3d three times over and select_2d_config against select_3d_config. Those are the shape the rule exists for. D9's body already records this lesson from the other end: it was a name list, it missed three real defects, and it is now an AST pass with no names in it. Do not add the suffix list back. Python only. Every hit was Python, and splitting C bodies by brace depth is a heuristic that earned nothing -- it produced a span bug that matched a function's own `def` line inside another function. Evidence, not a finding: a variant legitimately diverges. What is handed over is the pair, so that the judgement gets made instead of skipped. Checked on ten PRs outside the corpus (batch 23): it fires on aiter#2478, "Fix GPU memory access fault in CK MoE FP4 kernel", pointing at `_flydsl_moe_sorting` still holding the allocation-size line that `_moe_sorting_impl` just had fixed. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * review-pr: run D11's worked example instead of believing it A sweep reported that D11's example was fiction: aiter#5220 adds two ints to `pa_sparse_prefill_kargs`, the rule body says three assertions in the same translation unit fail, and grepping the tree for `offsetof(` near that struct finds nothing. The header does hold twelve static_asserts and every one of them is tile or vector arithmetic, which made the conclusion look doubly confirmed. I rewrote the docstring and the rule body around it -- #5220 recast as the case D11 correctly stays silent on. Both greps read /mnt/raid0/zufa/aiter, a second checkout sitting on af81aa0bc from 2026-07-13. On the branch under review, `pa_sparse_prefill_opus_kernels.cu` asserts `sizeof(pa_sparse_prefill_kargs) == 112` and fixes every field offset through the PA_GFX1250_CO_ABI macro -- which is the second disguise, because a literal `offsetof(` does not appear either. Given the right root, `triage.py structabi` on #5220 prints PINNED-LAYOUT. The rule was right and the measurement was stale, so the rewrite is reverted and only the test survives. The test is the point: the example is now executed, not asserted in prose. One case pins that the struct really is pinned in this tree, one runs #5220's shape through the detector and requires PINNED-LAYOUT, one requires silence when the same diff updates the assertion. If the tree moves, the failure says so and names which passage to rewrite. The tree is found from the skill's own location, so there is no second checkout to pick wrong. Also corrects sibling_variants: 18.0% was measured against that same stale tree, and the merge target gives 18.3%. The conclusion does not move. Nothing else needs re-measuring -- dropped_parameter and untested_new_kernel read the diff and never the tree, and fetch.sh takes PROJECT_ROOT from `git rev-parse --show-toplevel`, so every artifact under /tmp/rev was always computed against the right root. Only the hand-run measurements were wrong. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * review-pr: say which tree a forensic just read symbols, twins, citest, structabi and siblings all answer questions about the merge target -- is this struct pinned, does a CI job scan this path, does a variant still carry this line. Pointed at a second checkout they answer about that one instead, in the same shape, with nothing to show the reader which repository was consulted. That cost a rule last week. A sweep read /mnt/raid0/zufa/aiter, 53 days behind, found nothing pinning `pa_sparse_prefill_kargs`, and reported D11's worked example as fiction. The follow-up confirmed it against the same checkout -- twice over, since the offsets are pinned through a macro and a literal `offsetof(` is absent from both trees -- and the rule body was rewritten around the mistake before a test caught it. So the modes print, on stderr, when the root they were handed is not the tree this file ships from, with both HEADs and their dates beside each other. Silent in normal use: fetch.sh derives PROJECT_ROOT from `git rev-parse --show-toplevel`, so it always passes the tree the skill lives in. It is the measurement run by hand that picks the wrong one. stderr specifically, and there is a test for that: fetch.sh tees these modes into $WORK/*.txt, and a warning on stdout would be filed as evidence. Another walks the CLI block and fails if a root-taking mode reads argv without the guard, because the one that skips it is the one that goes wrong quietly. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * review-pr: refuse to review a diff that was never fetched `gh pr diff > $WORK/pr.diff` was unchecked. When it fails the file is empty and the run carries on: the deriver answers `underivable` and emits all 52 rules, which is the intended fail-open and does not come back looking clean -- it comes back as 52 rules that cannot be answered against a diff that is not there, while the reason sits on a stderr line a batch run discards. The reason also names the wrong cause. GitHub refuses a diff over 20000 lines with `HTTP 406: could not find pull request diff`, which reads as a bad PR number. aiter#4961 is one of the 240 open PRs scanned so far and is exactly this: a live cross-repo PR, 27k lines, reported as not found. An eighth guard, in the shape of the other seven. Over the cap it says so, says it is not a missing PR, and says to review from a local checkout of the head ref instead. Any other failure quotes gh's own stderr line by line rather than replacing it with a summary. 0.4% of PRs scanned, which is below the rate at which detection candidates have been turned down here -- and the rate is the wrong test for it. A detector at 0.4% finds almost nothing; this is not finding anything, it is refusing to proceed on nothing, and it has no false-positive side. Verified on aiter#4961, which now exits 1 with the cap named, and on aiter#5120, which still fetches and derives 15 rules. Dropping the emptiness check turns one of the four new tests red. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * review-pr: make T5 and T6 fire on the shape, not on the language Both are red rules and both were triggered by tokens that say "this is Triton code". `tl.cdiv(` was one of T6's triggers. It is ceiling division: of the seven firings it produced on its own across the corpus's 232 Triton PRs, all seven were a tile count, a block-pointer `shape=`, a mask bound or a loop bound. None computed a grid. T6 asks whether the host grid disagrees with the kernel's program_id axes, and a diff that merely divides something upward is not being asked anything. 27% of Triton PRs to 24%. A bare `.to(tl.float32)` was one of T5's. It produced 20 of that family's 54 firings and every one sampled was an upcast on a load, a store, or a quantisation max -- `m = tl.max(tl.abs(x)).to(tl.float32)` -- which is the correct direction, the opposite of the accumulator left in fp16 that T5 is about. 23% to 19%. Naming the accumulator instead is not only subtraction. `acc +=` and `acc = tl.zeros` bring in aiter#3613, `acc_sq += gl.sum(out_h * out_h, axis=1)` in a Gluon kernel with no `tl.dot` anywhere in the diff -- T5's first failure shape exactly, and invisible to the old trigger. The blessed corpus moves on two PRs: 3613 gains the family, 4268 loses it. Not landed, with the numbers, so the next pass does not repeat it: a forensic for T6 that reads both halves out of the tree and compares grid arity against the highest program_id axis. Four versions, mismatch rates of 50%, 40%, 34%, each drop a parsing bug of mine rather than a real narrowing. What ended it was that the same kernel and grid recurred across seven unrelated PRs -- so the "mismatch" was a property of the tree, and a real one would be a live bug on main that everyone hits. The cause is that resolving `kernel[grid](` means finding which `grid` binding is live at that call site, and a line-based search finds the first one in the file. That needs a scope-aware pass, which is what D9's scanner had to become for the same reason. The trigger-precision finding above does not depend on any of it: only 46% of T6's firings carry both halves in the diff at all, and only 50% of T5's carry a dot. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * review-pr: say what became of each guard the diff deletes D4 is red, `invariant-removed` fires on 69 of 600 open PRs, and the only forensic behind it collects symbols from deleted guards and greps for them. Its own docstring says the usual answer is "it moved" -- which the reviewer then has to establish by hand, on every one of the 69. This started as a filter and the numbers supported one: 14.5% of those PRs add every deleted guard straight back verbatim, and another quarter return them in a form that shares the subject, so about 40% looked like it could be dropped before the reviewer ever saw it. Reading seven of that second group is what changed the design. aiter#4295 replaces `AITER_CHECK(valid_split_count != nullptr && ...->data_ptr() != nullptr)` with a size check, and the null guard is simply gone. aiter#4279 turns `TORCH_CHECK(hidden_size >= (tile_k * split_k) * 2)` into the same check without the `* 2`. aiter#5168 widens `assert d_qk_v == (...)` to a set membership. Three of seven were real weakenings. "The check came back" and "the check was weakened" produce identical evidence under any same-subject test, and the second is the finding, so filtering on it would have deleted D4's best cases as noise. So it prints the pair instead. Each deleted guard is sorted into moved unchanged, returned changed, or gone, and the changed ones come out as before/after lines. Across the corpus: 14.5% are entirely moves and have nothing to review, 58% drop at least one guard outright, 46% return at least one changed. Two of the tests are aiter#4295's and aiter#4279's real lines, so the case for printing rather than filtering is pinned where the next person tempted to filter will run into it. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * review-pr: run the symbol sweep's worked example too sweep_symbols explains itself with aiter#4994: it added `from aiter.ops.flydsl.utils import is_flydsl_available`, #5116 had deleted that module from main 19 hours earlier, and #4994 merged green and was reverted 5.5 hours later. Whether the check fires at all depends on which tree it resolves against, which is precisely what prose cannot hold onto -- D11's example was prose in the same way, and reading it against a checkout two months stale produced a confident wrong conclusion that the rule was fiction. Four cases now: the named module is still absent from this tree, a module that exists is not reported, a module the diff itself adds is not reported, and a third-party import is left alone. The second is the control; without it the first passes against a sweep that flags everything. No gap this round. Replaying the sweep over the corpus reproduces the population it was built for: 61 of 600 open PRs import something unresolvable, 57 at module level across 26 modules, `aiter.ops.flydsl.utils` in 32 of them and `aiter.ops.triton.utils.core` in 7. Recall against that set is complete. Measured and turned down: the reverse direction, a PR deleting a module the tree still imports. 3 of 600 raw, and all three are the same false positive -- the PR deletes the importing line as well, which aiter#5145 shows on its face with `-from aiter.ops.flydsl.kernels.hgemm_dispatch import compile_flydsl_hgemm_kernel`. 0 of 600 real. The cross-PR case that actually bites -- one PR deleting a module 32 open PRs import -- needs the open-PR set, which fetch.sh does not have and would cost 600 diffs to get; the symptom side is already covered on each of those 32 when they are reviewed. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * review-pr: FlyDSL is a kernel backend, so derive the kernel rules for it KERNEL_PY lists `aiter/ops/triton/`, `_triton_kernels/`, `_gluon_kernels/` and `/gluon/`. A FlyDSL kernel matches none of them, so it reached `flydsl` -- D10 and D10b, which are about what happens to a compile result -- and nothing else. 108 of 600 open PRs edit `aiter/ops/flydsl/kernels/*.py`; 95 of them derive neither kernel family. A1's sibling variant, D1's uninitialised accumulator, D8's missing contiguous check and P6's unmeasured cost were never put to an entire backend. The siblings forensic had been pairing `_flydsl_stage1_wrapper` against `_flydsl_stage2_wrapper` in exactly these files for three commits. The evidence was being produced and no rule was asking for it. A separate family rather than another entry in KERNEL_PY, because FlyDSL is not Triton and `modified-kernel` would say it is. B2 is deliberately not in the set: it is `tl.load` or `tl.store` without a mask, and there is no tl here. Median rules per PR stays 15 and the mean moves 14.7 to 15.1 -- these PRs already carried most of what this adds. Eight of the pinned 59 gain the family and none loses anything. Found while auditing something else, which is worth recording: `perf` fires on a title keyword, and `fuse` accounts for 43 of its 90 firings on its own. Every one of twelve read was a kernel NAME -- "Expose the fused_moe activation dtype resolution", "Add `repr` to DiT fused kernels" -- not a claim about speed. Dropping `fuse` and `%` takes the family from 15% to 7.5%, and the reason that is not in this commit is that it interacts with this one: some of what it drops was reaching P6 only through `perf`, and now reaches it through `flydsl-kernel` instead. It needs re-measuring against this baseline, not the old one. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * validate-kernel-pr: say when another validator holds the GPU Two of these suites at once make a third of the run look broken. The validator claims a GPU with `flock -n` on /tmp/gpu-N.lock; a second one finds every candidate locked, degrades to NO_GPU, skips the runtime stages and returns INCONCLUSIVE. That is correct, and the report says so -- "GPU claim raced with another process". What the reader sees is `'PASS' != 'INCONCLUSIVE'` on five or six tests, which names the symptom and hides the cause one file away. Run alone: 89 passed, six separate times. Two runs overlapping: 5 failures and 6. Both of the overlaps were mine -- a run started with a bare `&`, then a second launched before it finished, twice. I spent several sessions reporting this suite as flaky and proposing to isolate the timing-sensitive tests, which would have deleted a true signal from a suite that was working. The reason never reached me because I had been piping the output through `tail -3`, so the one line that explained it was cut off every time. The guard asks the lock rather than the process list: it tries to flock each /tmp/gpu-*.lock and reports the ones it cannot take. Written first as `pgrep -af validate_pr.sh`, which matched the shell running the pgrep -- the string is in its own argv. Verified by holding a lock and watching it fail with the right message. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * review-pr: fix what four end-to-end runs found Everything until now was measured against the 600-PR corpus, which exercises the deriver and the forensics and never writes a card. Four full Step 1-8 reviews -- aiter#4295 and #2478 for signal, #3836 and #5072 as controls -- found six defects in one pass. Both controls returned NO FINDINGS, so the gates are not manufacturing work; what they were doing was corrupting the reviews that had something to say. The citation gates punished the forensics the rules ask for. card and ledger rejected any finding naming a file the diff does not change, and three of the four runs hit it independently. aiter#2478's two strongest findings rest on the mask contract in csrc/include/moe_sorting_opus.h -- an unchanged header, and the only place the convention and the local-id derivation are written down -- so the review deleted the filename and described the header in prose to get through. E4's own rule body sends the reader to .github/workflows/*.yaml to check a label, and doing so tripped the same gate. Both now require an anchor: at least one cited file the PR changes, with everything else free. A finding citing nothing from the diff is still rejected, which is what these were for. The card gate demanded rule codes that SKILL.md forbids. It matched a finding to its FIRE by a leading `D8:`, and SKILL.md line 466 says "Do NOT use rule codes (P1, D4, A1...) in output -- they are internal labels only". Every card written to spec reported every fired rule as UNREPORTED-FIRE. A finding naming a file the verdict cited is the same claim without the label. `torch.cuda.device(...)` parsed as a file called torch.cu, and `x.contiguous.cuda()` as x.contiguous.cu. Those are the two most ordinary expressions in a HIP review. SKILL.md told the reviewer that a runtime change with no test target is a finding in its own right. Triage calls anything under aiter/ a runtime surface, so both control PRs -- a tuner input CSV and a tuned-config table, neither loaded at run time -- arrived carrying a defect the reviewer was instructed to report. That is a fabrication inducer on the family of PRs where the deriver has least to say. Two artifacts could not say they had run. An empty symbols.txt means "this axis was not checked" per SKILL.md, and a clean sweep wrote zero bytes; ci_coverage said "every new test file lands where a CI job will run it" on a PR that adds no test file, which reads as a pass. And the artifacts line promised evidence.txt unconditionally though it is written only for guard or signature diffs, while never learning about guards.txt, siblings.txt or kernel_tests.txt -- three files I added, three commits apart, without touching that line. Also drops two .pyc files that had been force-added into validate-kernel-pr. They are build output and would land in any upstream PR. Not fixed, recorded for the next pass: twins compares within a language, so aiter#4295 -- a .cu to .py port, which is twin divergence by definition -- got "no new file closely mirrors an existing one" while the review found three divergences by hand; perf_claims par…
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
1 participant
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.
Summary
Migrates four gfx950 attention/delta FlyDSL kernels off raw MLIR dialects onto the
fx.*surface (part of a broader raw-dialect cleanup, split by area). Op-surface only — no logic/tiling/math/offset changes.gdr_decodevector.*,arith.*chunk_gated_delta_hllvm.inttoptr<3>, redundant_to_rawpa_mqa_logits_fp4buffer_ops.*,ArithValuepa_mqa_logits_fp4_prefillbuffer_ops.*,ArithValueWhat changed
vector.*→fx.Vector(.filled/.from_elements/fx.math.fmawhich keeps the fusedv_fma_f32).buffer_ops.create_buffer_resource+buffer_load/store→make_buffer_tensor+ copy atoms / indexed fx pointers; 1-writer scatter → guardedfx.ptr_store.arith.constant/ArithValue→ fx literals/operators;llvm.inttoptr<3>→fx.to_llvm_ptr(backend-resolved address space).scf.forloop-carry alreadyrange(init=); removed the redundant boundary_to_raw(the rewriter converts fx atinit=/yield).Kept raw (evidenced boundaries)
rocdl.mfma*,rocdl.ds_read_tr16_b64,rocdl.ds_bpermute,rocdl.exp2,gpu.shuffle(hand-packed MFMA operands / warp shuffles — no fx wrapper), and_llvm.mlir_undef(fx has no undef/poison). Tiled-copy (make_tiled_copy/partition_S/D) does not apply to any of these — they're per-lane MFMA / swizzled-LDS / paged-gather patterns, not contiguous bulk copies.Validation
gfx950, cache ON. Byte-exact vs the original per kernel (several proven via md5-identical lowered IR / ISA diff — stronger than
torch.equal);pa_mqaop tests pass before+after; no perf regression. Ruff + Black clean.Note: the
_gfx1250/_gfx1201attention variants are intentionally not in this PR — they can't be byte-exact/perf validated on a gfx950 box and need a matching-arch machine.