Skip to content

[Kimi-K3] FlyDSL SiTUv2 MoE kernels, strided grouped-topk, tuned configs - #4397

Merged
valarLip merged 43 commits into
mainfrom
k3-for-amd
Jul 27, 2026
Merged

valarLip merged 43 commits into
mainfrom
k3-for-amd

Conversation

@carlushuang

Copy link
Copy Markdown
Collaborator

Summary

Brings up Kimi-K3 (KDA-hybrid, MXFP4 MoE) in aiter, covering both gfx950 (MI355) and gfx1250. All changes are additive and gated behind K3-specific paths (SiTUv2 activation, K3 shapes); default behavior for existing models is unchanged.

Changes

FlyDSL SiTUv2 MoE kernels (a16w4)

  • flydsl_moe_stage1/2 with SiTUv2 activation (situ_beta, situ_linear_beta): mixed 2-stage a16w4 MoE GEMM, fused route/quant/scatter, opt-sort and shuffle utilities for the MXFP4 (a16w4) path.
  • flydsl 0.2.4 compatibility: drop loc= from _split_mfma_operands calls, _MIN_FLYDSL_VERSION bumped to 0.2.4.
  • Observable tile-alignment guard: force tile_n/k 256→128 when inter_dim is not 256-aligned (e.g. 384), with a one-time warning.

A4W4 / A8W4 SiTUv2 (gfx950 + gfx1250)

  • FlyDSL A4W4 SiTUv2 path and AITER_SITUV2_A8W4 env-gated a8w4 mode; a8w4 fmoe tune configs.

Grouped topk: row-strided inputs

  • biased_grouped_topk and the opt-sort kernel accept row-strided logits (the K3 fused MoE-front router slice of the [gate_up|experts|routed] buffer) — avoids a per-layer contiguous copy. Coverage folded into the existing test_biased_grouped_topk as a K3-shaped case (no new test file).

Tuned configs (Kimi-K3)

  • kimik3_fp4_tuned_fmoe.csv, kimik3_fp4_untuned_fmoe.csv, kimik3_bf16_tuned_gemm.csv, kimik3_bf16_untuned_gemm.csv, a16w4 fmoe tuned config (inter_dim=384), a8w4 fmoe tune configs, bf16 dense GEMM tuned config, gfx1250 conv2d (ViT) configs, gfx1250 bf16 tuned config.

Vision / misc

  • Triton conv2d configs for the K3 ViT tower (gfx1250), Gluon reduce guard, MlaVersion enum in module_aiter_core, CK gfx1250 enable patch (temporary).

billishyahao and others added 30 commits July 8, 2026 16:48
* add Situv2 activation for a8w4 MoE stage1

* refactor(flydsl): scope situv2 helpers and fix lint

Move situ beta compile-time constants into situ_elem/situ_up_elem; apply
ruff/black fixes on the a8w4 situv2 test module.

* feat(moe): integrate SiTUv2 into fused_moe API, split-K, tuner/prebuild

* test(moe): add SiTUv2 host-ref + a4w4/a8w4 stage1 test (adapted to Situv2 naming)

Complements MHYang's aiter/ops/flydsl/test_flydsl_moe_a8w4.py by adding a
host-only (no-GPU) SiTUv2 reference sweep plus a4w4 stage1 coverage in
addition to a8w4. Adapted to this branch's API: ActivationType.Situv2 enum
and situ_beta / situ_linear_beta parameters on torch_moe_stage1 and
flydsl_moe_stage1.

Co-authored-by: Cursor <cursoragent@cursor.com>

* test(moe): add SiTUv2 default cases + --beta/--linear-beta to test_moe_2stage (Situv2 naming)

Co-authored-by: Cursor <cursoragent@cursor.com>

* test: consolidate a8w4 SiTUv2 vec4 cases into op_tests/flydsl_tests.

Move pytest tile/gate_mode sweep from aiter/ops/flydsl/test_flydsl_moe_a8w4.py
into test_flydsl_moe_situv2.py and remove the misplaced test file.

* style: black-format silu_and_mul_fq.py for CI pre-checks.

---------

Co-authored-by: MHYang <mengyang@amd.com>
Co-authored-by: Clement Lin <Clement.Lin@amd.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
* add Situv2 activation for a8w4 MoE stage1

* refactor(flydsl): scope situv2 helpers and fix lint

Move situ beta compile-time constants into situ_elem/situ_up_elem; apply
ruff/black fixes on the a8w4 situv2 test module.

* feat(moe): integrate SiTUv2 into fused_moe API, split-K, tuner/prebuild

* test(moe): add SiTUv2 host-ref + a4w4/a8w4 stage1 test (adapted to Situv2 naming)

Complements MHYang's aiter/ops/flydsl/test_flydsl_moe_a8w4.py by adding a
host-only (no-GPU) SiTUv2 reference sweep plus a4w4 stage1 coverage in
addition to a8w4. Adapted to this branch's API: ActivationType.Situv2 enum
and situ_beta / situ_linear_beta parameters on torch_moe_stage1 and
flydsl_moe_stage1.

Co-authored-by: Cursor <cursoragent@cursor.com>

* test(moe): add SiTUv2 default cases + --beta/--linear-beta to test_moe_2stage (Situv2 naming)

Co-authored-by: Cursor <cursoragent@cursor.com>

* test: consolidate a8w4 SiTUv2 vec4 cases into op_tests/flydsl_tests.

Move pytest tile/gate_mode sweep from aiter/ops/flydsl/test_flydsl_moe_a8w4.py
into test_flydsl_moe_situv2.py and remove the misplaced test file.

* style: black-format silu_and_mul_fq.py for CI pre-checks.

* Fix GUI shuffle_scale k_pad for non-256-aligned MoE K and auto stage2 tile_k.

Pad w2/w1 GUI e8m0 scales to k_groups multiple of 8 (DSV4 inter=640), unify
stage2 tile_k selection in flydsl_moe_stage2 and fused_moe, and add staged
a8w4 regression tests across inter/model K sweeps.

* test: integrate a8w4 K-dimension regressions into existing test suites.

Move shuffle pad coverage into test_quant_mxfp4 and FlyDSL GUI stage2/e2e
into test_flydsl_moe_a8w4; drop standalone op_tests files for repo convention.

* test: move flydsl a8w4 GUI regressions to op_tests/flydsl_tests.

Relocate test_flydsl_moe_a8w4 under op_tests/flydsl_tests and align with
pytest conventions used by other FlyDSL op tests.

---------

Co-authored-by: MHYang <mengyang@amd.com>
Co-authored-by: Clement Lin <Clement.Lin@amd.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
* add Situv2 activation for a8w4 MoE stage1

* refactor(flydsl): scope situv2 helpers and fix lint

Move situ beta compile-time constants into situ_elem/situ_up_elem; apply
ruff/black fixes on the a8w4 situv2 test module.

* feat(moe): integrate SiTUv2 into fused_moe API, split-K, tuner/prebuild

* test(moe): add SiTUv2 host-ref + a4w4/a8w4 stage1 test (adapted to Situv2 naming)

Complements MHYang's aiter/ops/flydsl/test_flydsl_moe_a8w4.py by adding a
host-only (no-GPU) SiTUv2 reference sweep plus a4w4 stage1 coverage in
addition to a8w4. Adapted to this branch's API: ActivationType.Situv2 enum
and situ_beta / situ_linear_beta parameters on torch_moe_stage1 and
flydsl_moe_stage1.

Co-authored-by: Cursor <cursoragent@cursor.com>

* test(moe): add SiTUv2 default cases + --beta/--linear-beta to test_moe_2stage (Situv2 naming)

Co-authored-by: Cursor <cursoragent@cursor.com>

* test: consolidate a8w4 SiTUv2 vec4 cases into op_tests/flydsl_tests.

Move pytest tile/gate_mode sweep from aiter/ops/flydsl/test_flydsl_moe_a8w4.py
into test_flydsl_moe_situv2.py and remove the misplaced test file.

* style: black-format silu_and_mul_fq.py for CI pre-checks.

* Fix GUI shuffle_scale k_pad for non-256-aligned MoE K and auto stage2 tile_k.

Pad w2/w1 GUI e8m0 scales to k_groups multiple of 8 (DSV4 inter=640), unify
stage2 tile_k selection in flydsl_moe_stage2 and fused_moe, and add staged
a8w4 regression tests across inter/model K sweeps.

* test: integrate a8w4 K-dimension regressions into existing test suites.

Move shuffle pad coverage into test_quant_mxfp4 and FlyDSL GUI stage2/e2e
into test_flydsl_moe_a8w4; drop standalone op_tests files for repo convention.

* test: move flydsl a8w4 GUI regressions to op_tests/flydsl_tests.

Relocate test_flydsl_moe_a8w4 under op_tests/flydsl_tests and align with
pytest conventions used by other FlyDSL op tests.

* Add FlyDSL MoE a16wfp4 (bf16 x mxfp4) kernels on gfx950.

Port a16w4 stage1/stage2 into mixed_moe_gemm_2stage as dedicated _a16w4
builders while keeping HEAD fp8/fp4 compile paths unchanged. Wire dispatch
in moe_kernels and add op_tests/flydsl_tests regression for stage1, stage2,
and e2e against torch_moe references.

* Fix ruff/black lint issues in a16wfp4 PR files.

Remove unused _fly imports, move lds_space/types imports to module top,
add noqa for test sys.path bootstrap, and apply black formatting.

* Remove dead duplicate waves_per_eu None checks in a16w4 stage2.

The parameter defaults to int and callers always pass an int after dispatch.

* Fix a16w4 stage2: skip second compute when total_tiles==1.

The CK-style ping-pong pipeline assumed total_tiles>=2.  When
inter_dim==tile_k (e.g. inter_dim=256, tile_k=256) total_tiles=1,
the HEAD pre-fetched an out-of-bounds A tile into lds_ping, and the
TAIL then computed with that garbage data, producing cos~0.71 vs
reference.

Fix: const_expr(total_tiles==1) guard in the TAIL block--skip the
second compute_tile and the ping-buffer load entirely.

Also tighten test _check_result: add cosine+rel_L2 as primary gate
(cos>0.999) so future stage2 bugs cannot be masked by atol=1.0 on
small-magnitude outputs.

* Add pipeline guards for a16w4 stage1 total_tiles==1 and odd-tile validation.

The stage1 ping-pong pipeline has the same total_tiles>=2 assumption
as stage2.  Add:
- total_tiles==1 guard in stage1 TAIL (same pattern as stage2 fix)
- Validation: reject odd total_tiles and K < tile_k at compile time
  for both stage1 and stage2 a16w4 paths

These configs cannot occur with real LLM shapes (model_dim and
inter_dim are always multiples of 256), but the guards prevent
silent corruption on contrived inputs.

* Add SiTUv2 activation to a16w4 (bf16 x mxfp4) stage1.

Port the SiTUv2 activation from pr/a8w4-situv2 into the a16w4 stage1
kernel (both direct and cshuffle epilogues):
  situ_g   = beta * tanh(gate/beta) * sigmoid(gate)
  up_scaled = linear_beta * tanh(up/linear_beta)
with gate/up clamped to <=7 / [-7,7] before activation, matching the
a8w4 situv2 numerics exactly.

Thread situ_beta/situ_linear_beta through compile_flydsl_moe_stage1 and
flydsl_moe_stage1; the per-beta cache tag keeps distinct binaries.

Verified vs an inline torch SiTUv2 reference (mxfp4 dequant GEMM +
clamp + situv2): cos>=0.99996 across beta in {(1,1),(0.5,2),(1.5,0.8)},
shapes (512/256, 3072/256), tile_n {128,256}. silu/swiglu unchanged.

* Strip dead generic pipeline from a16w4 stage1 kernel (~2260 lines).

compile_mixed_moe_gemm1_a16w4 was created by copying the generic fp8/fp4
builder and gating with is_a16w4_stage1.  The live a16w4 path is fully
contained in `if const_expr(is_a16w4_stage1): ... return`; everything
after that return was an unreachable copy of the generic f8f6f4 kernel
body (referenced generic-only vars like _lds_tid_offset_pong / _pipe_*
that are never defined on the a16w4 path -- proven dead since a16w4
tests pass without them).

Remove the dead generic kernel body, the generic-only setup blocks
(postlude pipeline schedule, not-is_a16w4 guards), and the now-unused
setup vars.  Purely dead-code deletion: a16w4 stage1/stage2/e2e/situv2
+ a4w4 all still pass with identical cosine.

* Strip dead generic pipeline from a16w4 stage2 kernel (~1140 lines).

Same shape as the stage1 cleanup: compile_mixed_moe_gemm2_a16w4's live
a16w4 path is fully contained in `if const_expr(is_a16w4): ... return`
inside _moe_gemm2_then_body; everything after that return was an
unreachable copy of the generic f8f6f4 stage2 body (f8f6f4 MFMA,
generic-only scale layouts). Proven dead: a16w4 stage2/e2e/situv2 +
a4w4 all pass unchanged after removal.

Remove the dead body plus the now-unused generic setup vars
(pack_N/pack_K, cbsz/blgp, generic scale layouts, sx_rsrc sentinel,
etc.). Purely dead-code deletion.

* feat(moe): route a16w4 (bf16 x mxfp4) SiTUv2 through fused_moe

Wire the mixed_moe a16w4 kernel into the fused_moe 2-stage path for SiTUv2:
get_2stage_cfgs now matches bf16/fp16 x fp4 when activation is SiTUv2 (which
uniquely identifies the a16w4 kernel, so GPT-OSS / legacy bf16-Swiglu keep
their CK-Tile routing), maps _a_type to bf16/fp16, and infers q_dtype_a=bf16
for SiTUv2+separated so the activation stays bf16 (no fp4 quant). The stage1/
stage2 activation-quant gates also accept SiTUv2.

Tests: a16w4 SiTUv2 via fused_moe (cos=1.0 vs torch ref) and an a8w4 SiTUv2
inter=640 vec4 case exercising the fix-k non-256 K-tiling.

* feat(moe): support non-256 inter_dim for a16w4 (a8w4 parity)

Bring the FlyDSL a16w4 (bf16 x mxfp4) MoE kernels to a8w4 parity so
inter_dim need not be a multiple of 256 (e.g. DSV4 inter=640, and
arbitrary values like 384 via inter_dim_pad), across stage1/stage2/E2E.

Kernel (mixed_moe_gemm_2stage.py):
- stage2: allow inter_dim % tile_k == 0 (drop even-only rule), add
  odd_k_tiles 1-tile tail path, 256-pad scale reads (#3476).
- stage1: same odd_k_tiles K-loop + 256-padded w1 scale layout; fix gx
  launch grid (tile2_pad + full-N when inter_dim_pad > 0).

Caller (moe_kernels.py):
- pick/resolve_flydsl_stage1_tile_n: use tile_n=128 when inter%256!=0
  (fixes stage1 error on cols 0..inter%256 under tile_n=256).
- zero-init stage1 out when inter_dim_pad > 0.

Tests (test_flydsl_moe_a16wfp4.py):
- non-256 regressions (256/384/640) for stage1/stage2/E2E + pick_* tests.
- _generate_a16wfp4_data gains activation/situ_beta/situ_linear_beta.
- --perf sweep mode (correctness + latency/TFLOPs).

* test(moe): fold situv2 tests into dtype files, drop standalone situv2 file

Consolidate test_flydsl_moe_situv2.py into the per-dtype test files (matching
how test_flydsl_moe_a16wfp4.py already embeds its situv2 tests):

- test_flydsl_moe_a8w4.py: add the a8w4 SiTUv2 vec4 stage1 sweep + the
  host-only situv2 reference test; reuse the file's existing _check_close
  helper (fp32 cast for the bf16-ref vs f16-out case). Add shuffle_weight /
  e8m0_shuffle imports.
- Delete test_flydsl_moe_situv2.py. Its broken-under-pytest a4w4/a8w4
  main()-runner is dropped; a8w4 situv2 is now real parametrized coverage.

a16wfp4 situv2 tests already lived in their dtype file and are unchanged.

* fix conflict

* add timing args

---------

Co-authored-by: MHYang <mengyang@amd.com>
Co-authored-by: Clement Lin <Clement.Lin@amd.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: mh <mh@local>
…n tests

Squashed fix + tests (PR#3 commits 50729303, 2f33bd40, 1f0c6485, 46e3e056).

flydsl_moe_stage1 only downgraded tile_n (256->128) for non-256-aligned
inter_dim on a16w4 (bf16 x mxfp4); a8w4 (fp8 x mxfp4) kept tile_n=256, which
in separated gate_mode over-runs the gate/up (N) axis for non-256 inter_dim:
~30% wrong E2E output or GPU memfault at inter=384/640. Extend the resolve to
a8w4 (b_dtype fp4/mxfp4, a_dtype in {bf16,fp8}); a4w4 untouched. Callers keep
passing tile_n=256; the kernel resolves internally.

Tests (numeric vs torch reference, no-pad path, 128-multiple inter_dim):
- test_flydsl_e2e_a8w4_situv2 (separated, the production/customer path)
- test_flydsl_e2e_a16wfp4_situv2 (separated + interleave)
- test_flydsl_e2e_a8w4_gui extended to inter=256/384/640 (interleave, swiglu)

Verified on gfx950 (full rebuild): a8w4 20 / a16wfp4 27 / a4w4 4 passed;
ruff + black clean.
…tuv2 E2E

Squashed (PR#3 c827ddbb + d381aae7).

- resolve_flydsl_stage1_tile_n / resolve_flydsl_stage2_tile_k: keep the silent
  auto-downgrade of a non-dividing tile (256->128) for non-256 inter_dim, but
  make it observable -- full docstrings noting tile=256 is NOT tunable for such
  shapes, plus a one-time (deduped) logger.warning on override.
- test_flydsl_e2e_a8w4_situv2: enable interleave (a16w4-style shuffle_weight_a16w4
  recipe) alongside separated, over 128-multiple inter_dim. Verified all
  activations (silu/swiglu/situv2) x both gate_modes x non-256 give E2E 0.0000.
Needed by atom-k3 (rocm/atom HEAD imports them from aiter.ops.shuffle);
ported verbatim from aiter main. Keeps the single aiter-k3 branch usable
by vllm-k3, sglang-k3 and atom-k3.
Bring k3-for-amd up to rocm/aiter main (6c0c6f6, +25 commits) on top of the
K3 kernel work (fda97e9 shuffle ops).

One conflict in aiter/ops/shuffle.py (GUI-interleave scale padding): rocm's
version kept — its if k_%8!=0 block already pads with proper fill, making the
K3 block redundant (identical padded shape; scale fill value is irrelevant
since the padded weights are zero). interleave_gate_up_rows +
moe_shuffle_weight preserved.
[kimi-K3] extend conv2d support to gfx1250
flydsl 0.2.4 removed the loc= parameter from rocdl._split_mfma_operands.
The A16W4 BF16 K32 MFMA helper in mixed_moe_gemm_2stage.py called it as
_split_mfma(operands, loc=loc), which raises TypeError under 0.2.4.

Drop loc= from both (identical) helper sites; _split_mfma_operands only
unwraps operand Values and reads int flags, so it never needed loc. The
loc/ip are still threaded to the real op builder (_mfma_k32_raw), so
location tracking is preserved.

Bump _MIN_FLYDSL_VERSION to 0.2.4 and drop the now-stale 0.2.2 pin note.

Verified on gfx950 with flydsl 0.2.4 via test_flydsl_moe_a16wfp4.py
(stage1, stage2, e2e all pass, cos>=0.99999).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
[Kimi-K3] fix MoE A16W4 for flydsl 0.2.4 and bump _MIN_FLYDSL_VERSION
[Kimi-K3] extend a16w4 fmoe tuned config inter_dim=384
Keep Kimi-K3 SiTUv2 on the A4W4 FlyDSL path, add gfx1250 grouped-MoE activation and stability support, preserve ATOM's legacy A4W4 call, and cover the gfx1250 path with focused tests.

Co-authored-by: Cursor <cursoragent@cursor.com>
Signed-off-by: Hongxia Yang <hongxia.yang@amd.com>
Enable a8w4 SiTUv2 MoE via AITER_SITUV2_A8W4
patches/ck_gfx1250_enable.patch only served to build native CK kernels
(quant/cache/rmsnorm/moe/sample/custom_all_reduce) on gfx1250. Verified that
with ENABLE_CK=0 (no patch, triton/hip fallbacks) Kimi-K3 serves correctly on
gfx1250 x4 (tp4): full gsm8k 1319 = 0.956 flexible/strict, vs 0.9613 with the
patched native-CK path — within stderr (+/-0.0056). Dropping the patch keeps
cleanup/k3-minimal minimal for merge into k3-for-amd; gfx1250 runs ENABLE_CK=0
(cost: ~13% decode vs native CK, no accuracy loss).
Kimi-K3 runs full-attn via MLAAttention (MLA-latent) and KDA via fla, so it never
dispatches to unified_attention (verified 0 calls across a full gsm8k run). The
non-power-of-two head_size / Gluon-reduce guards were added for K3's old
MHA-via-unified_attention path, which ATOM has now removed. Reverting
unified_attention.py to k3-for-amd keeps cleanup/k3-minimal focused on what K3
actually needs. Full gsm8k 1319 with this + the dropped CK patch = 0.955.
[Kimi-K3] add a8w4 fmoe tune config
Resolve aiter/fused_moe.py SiTUv2 conflict by taking k3-for-amd's a16w4/a8w4
logic: on gfx1250 q_dtype_a is unconditionally overridden below (fp4x2, or a8w4
via AITER_FORCE_A8W4), so Kimi-K3 is unaffected by this branch; the upstream
version is more correct for non-gfx1250. PR#14's remaining slimming (drop CK
gfx1250 patch, drop unified_attention Gluon guard) is preserved.
@carlushuang
carlushuang requested a review from a team July 27, 2026 15:12
@github-actions

Copy link
Copy Markdown
Contributor

🏷️ CI Guide

Runs automatically on every PR:

  • ✅ Pre-checks (submodule verification, code formatting)
  • ✅ Aiter op tests (gfx942 + gfx950)
  • ✅ Triton tests on MI35X (only when aiter/ops/triton/** or related paths are changed)

Extended tests (opt-in via labels):

Label Tests
ci:triton-300x Run an additional Triton test job on MI300X in PRs; main branch always runs both MI35X and MI300X
ci:sglang SGLang integration tests: DeepSeek-R1-MXFP4 accuracy, Qwen 3.5 accuracy
ci:atom ATOM benchmark: DeepSeek-R1-0528, GPT-OSS-120B
ci:atom_full ATOM accuracy suite for PR and main models from ATOM models_accuracy.json
ci:vllm vLLM benchmark: GPT-OSS-120B, DeepSeek-R1-0528, Kimi-K2.5
ci:all All standard extended tests (excludes ci:atom_full)

Only add ci:atom_full for FlyDSL or Triton upgrades.
Add labels via the sidebar or gh pr edit 4397 --add-label <label>

@github-actions github-actions Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Remaining comments which cannot be posted as a review comment to avoid GitHub Rate Limit

ruff

⚠️ [ruff] <BLE001> reported by reviewdog 🐶
Do not catch blind exception: Exception

except Exception: # pragma: no cover - logging must never break the kernel


⚠️ [ruff] <UP045> reported by reviewdog 🐶
Use X | None for type annotations

def get_flydsl_kernel_params(name: str) -> Optional[Dict]:


⚠️ [ruff] <UP006> reported by reviewdog 🐶
Use dict instead of Dict for type annotation

def get_flydsl_kernel_params(name: str) -> Optional[Dict]:


⚠️ [ruff] <UP045> reported by reviewdog 🐶
Use X | None for type annotations

w1_scale: Optional[torch.Tensor] = None,


⚠️ [ruff] <UP045> reported by reviewdog 🐶
Use X | None for type annotations

a1_scale: Optional[torch.Tensor] = None,


⚠️ [ruff] <UP045> reported by reviewdog 🐶
Use X | None for type annotations

sorted_weights: Optional[torch.Tensor] = None,


⚠️ [ruff] <UP045> reported by reviewdog 🐶
Use X | None for type annotations

k_batch_intra_block: Optional[int] = None,


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import pytest # noqa: E402
import torch # noqa: E402
import aiter # noqa: E402
from aiter import dtypes, QuantType, ActivationType # noqa: E402
from aiter.fused_moe import ( # noqa: E402
fused_topk,
moe_sorting,
torch_moe_stage1,
torch_moe_stage2,
)
from aiter.jit.utils.chip_info import get_gfx # noqa: E402
from aiter.ops.flydsl.utils import is_flydsl_available # noqa: E402
from aiter.ops.flydsl.moe_kernels import ( # noqa: E402
pick_flydsl_stage1_tile_n,
pick_flydsl_stage2_tile_k,
)
from aiter.ops.shuffle import shuffle_scale_a16w4, shuffle_weight_a16w4 # noqa: E402
from aiter.ops.quant import ( # noqa: E402
mxfp4_moe_sort_fwd,
per_1x32_f8_scale_f8_quant,
)
from aiter.utility.fp4_utils import e8m0_shuffle # noqa: E402


⚠️ [ruff] <RUF100> reported by reviewdog 🐶
Unused noqa directive (non-enabled: E402)

import pytest # noqa: E402


⚠️ [ruff] <RUF100> reported by reviewdog 🐶
Unused noqa directive (non-enabled: E402)


⚠️ [ruff] <RUF100> reported by reviewdog 🐶
Unused noqa directive (non-enabled: E402)


⚠️ [ruff] <RUF100> reported by reviewdog 🐶
Unused noqa directive (non-enabled: E402)

from aiter import dtypes, QuantType, ActivationType # noqa: E402


⚠️ [ruff] <RUF100> reported by reviewdog 🐶
Unused noqa directive (non-enabled: E402)

from aiter.fused_moe import ( # noqa: E402


⚠️ [ruff] <RUF100> reported by reviewdog 🐶
Unused noqa directive (non-enabled: E402)

from aiter.jit.utils.chip_info import get_gfx # noqa: E402


⚠️ [ruff] <RUF100> reported by reviewdog 🐶
Unused noqa directive (non-enabled: E402)

from aiter.ops.flydsl.utils import is_flydsl_available # noqa: E402


⚠️ [ruff] <RUF100> reported by reviewdog 🐶
Unused noqa directive (non-enabled: E402)

from aiter.ops.flydsl.moe_kernels import ( # noqa: E402


⚠️ [ruff] <RUF100> reported by reviewdog 🐶
Unused noqa directive (non-enabled: E402)

from aiter.ops.shuffle import shuffle_scale_a16w4, shuffle_weight_a16w4 # noqa: E402


⚠️ [ruff] <RUF100> reported by reviewdog 🐶
Unused noqa directive (non-enabled: E402)

from aiter.ops.quant import ( # noqa: E402


⚠️ [ruff] <RUF100> reported by reviewdog 🐶
Unused noqa directive (non-enabled: E402)

from aiter.utility.fp4_utils import e8m0_shuffle # noqa: E402


⚠️ [ruff] <C408> reported by reviewdog 🐶
Unnecessary dict() call (rewrite as a literal)

_s1_extra = dict(situ_beta=situ_beta, situ_linear_beta=situ_linear_beta)


⚠️ [ruff] <C408> reported by reviewdog 🐶
Unnecessary dict() call (rewrite as a literal)

return dict(
inter_pad=inter_pad,
ref_stage1=ref1,
ref_stage2=ref2,
inp=inp,
a2=a2,
w1_qt=w1_qt,
w1_scale=w1_scale,
topk_ids=topk_ids,
w1_qt_shuf=w1_qt_shuf,
w1_scale_shuf=w1_scale_shuf,
w2_qt_shuf=w2_qt_shuf,
w2_scale_shuf=w2_scale_shuf,
sorted_ids=sorted_ids,
sorted_weights_s1=sorted_weights_s1,
sorted_weights_s2=sorted_weights_s2,
sorted_expert_ids=sorted_expert_ids,
num_valid_ids=num_valid_ids,
dtype=dtype,
)


⚠️ [ruff] <C408> reported by reviewdog 🐶
Unnecessary dict() call (rewrite as a literal)

return dict(use_cuda_event=True)


⚠️ [ruff] <C408> reported by reviewdog 🐶
Unnecessary dict() call (rewrite as a literal)

return dict(use_cuda_event=False, testGraph=True)


⚠️ [ruff] <C408> reported by reviewdog 🐶
Unnecessary dict() call (rewrite as a literal)

return dict(use_cuda_event=False) # device


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

from aiter.test_common import run_perftest
from aiter.ops.flydsl.moe_kernels import flydsl_moe_stage1, flydsl_moe_stage2


⚠️ [ruff] <BLE001> reported by reviewdog 🐶
Do not catch blind exception: Exception


⚠️ [ruff] <BLE001> reported by reviewdog 🐶
Do not catch blind exception: Exception


⚠️ [ruff] <BLE001> reported by reviewdog 🐶
Do not catch blind exception: Exception


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

from __future__ import annotations
import os
import pytest
import torch
from aiter import dtypes, QuantType, ActivationType
from aiter.fused_moe import fused_topk, moe_sorting, torch_moe_stage1, torch_moe_stage2
from aiter.jit.utils.chip_info import get_gfx
from aiter.ops.flydsl.moe_kernels import (
pick_flydsl_stage2_tile_k,
resolve_flydsl_stage2_tile_k,
)
from aiter.ops.flydsl.utils import is_flydsl_available
from aiter.ops.quant import (
mxfp4_moe_sort_fwd,
per_1x32_f4_quant,
per_1x32_f8_scale_f8_quant,
)
from aiter.ops.shuffle import shuffle_weight, shuffle_weight_a16w4, shuffle_scale_a16w4
from aiter.utility.fp4_utils import e8m0_shuffle
from aiter.test_common import checkAllclose


⚠️ [ruff] <C408> reported by reviewdog 🐶
Unnecessary dict() call (rewrite as a literal)

return dict(
inter_pad=inter_pad,
topk=topk,
a_q=a_q,
a_scale_sort=a_scale_sort,
w1_shuf=w1_shuf,
w1_scale_shuf=w1_scale_shuf,
w2_shuf=w2_shuf,
w2_scale_shuf=w2_scale_shuf,
a2_q=a2_q,
a2_scale_sort=a2_scale_sort,
sorted_ids=sorted_ids,
sorted_weights=sorted_weights,
sorted_expert_ids=sorted_expert_ids,
num_valid_ids=num_valid_ids,
ref_stage1=ref_stage1,
ref_stage2=ref_stage2,
token=token,
inter_dim=inter_dim,
model_dim=model_dim,
)


⚠️ [ruff] <C408> reported by reviewdog 🐶
Unnecessary dict() call (rewrite as a literal)

return dict(
ref_stage1=ref_stage1,
a_q=a_q,
a_scale_sort=a_scale_sort,
w1_q_shuf=w1_q_shuf,
w1_scale_shuf=e8m0_shuffle(w1_scale),
sorted_ids=sorted_ids,
sorted_expert_ids=sorted_expert_ids,
num_valid_ids=num_valid_ids,
topk=topk,
)


⚠️ [ruff] <C408> reported by reviewdog 🐶
Unnecessary dict() call (rewrite as a literal)

return dict(
token=token,
inter_dim=inter_dim,
topk=topk,
a_q=a_q,
a_scale_sort=a_scale_sort,
w1_shuf=w1_shuf,
w1_scale_shuf=w1_scale_shuf,
w2_shuf=shuffle_weight_a16w4(w2_q, 16, False),
w2_scale_shuf=shuffle_scale_a16w4(w2_scale, E, False),
sorted_ids=sorted_ids,
sorted_weights=sorted_weights,
sorted_expert_ids=sorted_expert_ids,
num_valid_ids=num_valid_ids,
ref_stage1=ref1,
ref_stage2=ref2,
)


⚠️ [ruff] <C408> reported by reviewdog 🐶
Unnecessary dict() call (rewrite as a literal)

return dict(beta=args.beta, linear_beta=args.linear_beta)


⚠️ [ruff] <C408> reported by reviewdog 🐶
Unnecessary dict() call (rewrite as a literal)

yield dict(
dtype=dtype,
token=m,
model_dim=model_dim,
inter_dim=inter_dim,
E=args.expert,
topk=args.topk,
actType=aiter.ActivationType.Situv2,
gateMode=_effective_gate_mode(aq_dtype, wq_dtype),
qType=quant_type,
AQDType=aq_dtype,
WQDType=wq_dtype,
use_g1u1=True,
doweight_stage1=False,
strict_accuracy=False,
check_aot_cache=False,
swiglu_limit=None,
beta=args.beta,
linear_beta=args.linear_beta,
), extras

valarLip added 2 commits July 27, 2026 15:51
The pre-checks ruff job installs ruff unpinned (`pip3 install ruff`) and the
repo carries no ruff config, so the enforced rule set follows ruff's defaults.
ruff 0.16.0 widened those defaults well beyond E4/E7/E9/F, which is what this
branch tripped over. Fix the 41 findings reviewdog flagged inside the PR diff:

- silu_and_mul_fq.py (B023): bind the per-iteration SSA values as default args
  in _fmin / _sigmoid_s / _situv2_elem. The closures are deliberately redefined
  per unrolled iter_idx because the arith.constant ops must be emitted at the
  current insertion point, so hoisting them out of the loop is not an option.
- moe_kernels.py: narrow the logger-import guard to ImportError (BLE001) and
  switch the flagged annotations to PEP 604 / builtin generics (UP006, UP045).
- flydsl moe tests + test_moe_2stage.py: sort the import blocks (I001), drop
  the now-unused `# noqa: E402` markers (RUF100), rewrite dict() calls as
  literals (C408), and mark the sweep-runner blanket catches with an explicit
  `# noqa: BLE001` plus the reason they are intentional.

Only findings inside the PR diff are addressed; pre-existing findings elsewhere
in the same files are left alone. black --check stays clean.
reviewdog only posts a bounded number of results per run, so resolving the
first 41 findings surfaced a second batch of 27 that were in the PR diff all
along. Same root cause (unpinned ruff + no repo ruff config, 0.16.0 defaults).

- fused_moe.py / grouped_moe_gfx1250.py / moe_grouped_gemm_mxscale_gfx1250.py
  (UP045): PEP 604 annotations. Converted whole parameter runs rather than the
  individual flagged lines, because fixing one line pulls its neighbours into
  the diff hunk and the findings cascade.
- mfma_preshuffle_pipeline.py: drop the quotes from the ir.Value/ir.Type
  annotations (UP037) and sort __all__ (RUF022).
- mixed_moe_gemm_2stage.py: sort the import blocks (I001), functools.cache for
  the maxsize=None caches (UP033), inline the directly-called lambda in
  out_mlir (PLC3002), merge the nested bf16-atomics guard (SIM102), and
  underscore-prefix the unpacked values that are never read (RUF059) -- the
  decode calls are kept since they emit IR.

All rewrites are behaviour-preserving; black --check stays clean and the import
name sets are unchanged. Findings outside the PR diff are still left alone.
@valarLip
valarLip merged commit 3135022 into main Jul 27, 2026
24 checks passed
@valarLip
valarLip deleted the k3-for-amd branch July 27, 2026 16:05
@carlushuang
carlushuang restored the k3-for-amd branch July 27, 2026 16:07
@carlushuang

Copy link
Copy Markdown
Collaborator Author

valarLip added a commit that referenced this pull request Jul 28, 2026
test_gui_shuffle_scale_w2_k_pad has been failing on main at inter_tp=640:

    File "op_tests/test_quant_mxfp4.py", line 334
      assert torch.equal(
    AssertionError: inter_tp=640

The kernel is fine; the reference model in the test is wrong. Both sides
of the assert call the same shuffle_scale -- the only difference is who
pads k_groups up to a multiple of 8. shuffle_scale does it internally
with 0x7F for 1-byte scales (aiter/ops/shuffle.py:372), which in e8m0 is
the exponent bias, i.e. 2^0. The test pre-padded with torch.zeros, and
0x00 in e8m0 is 2^-127, so the two layouts genuinely differ.

inter_tp=256 gives k_groups=8, already a multiple of 8, so no padding
happens and the fill value never matters -- which is why only the 640
case failed.

Pad the reference with the neutral scale instead, mirroring what the
implementation guarantees: k-groups added by padding describe weights
that do not exist, so their scale has to be 1.0.

Verified on gfx950: with torch.zeros the assert fails for both float16
and bfloat16 at inter_tp=640 and passes at 256; with the neutral fill
all four combinations are byte-identical. `python3
op_tests/test_quant_mxfp4.py --a16w4_shuffle` goes from exit 1 to exit
0 on the same tree, and the full file passes.

Not a regression -- the test was written this way. The internal padding
and its 0x7F fill both landed in c2b69b9 (#4130, 2026-07-09);
test_gui_shuffle_scale_w2_k_pad was added later in 3135022 (#4397,
2026-07-28) and did not mirror the fill value.

Swept for the same assumption elsewhere: no other test hand-pads a scale
tensor before calling shuffle_scale. Every other caller passes unpadded
production data and relies on the internal padding, which is correct.
valarLip added a commit that referenced this pull request Jul 28, 2026
… sweep

Running op_tests/test_moeTopkSoftmax.py with its default token list dies
immediately:

    File "op_tests/test_moeTopkSoftmax.py", line 708, in <module>
      assert not gating_output.is_contiguous()
    AssertionError

The block builds a row-strided view of a fused router buffer:

    backing = torch.randn((token, 6016), dtype=bf16)
    gating_output = backing[:, 1536:1536 + 896]

and the default token list starts at 1. At token=1 the view is
(1, 896) with stride (6016, 1), and a row stride you can never take is
not a stride at all -- torch reports it as contiguous. The assert is
false for the first case it ever sees.

`assert gating_output.stride(0) == fused_front_width` is the check that
actually pins the layout under test, and it holds for every token. Keep
it and drop the other one.

Verified on gfx950: `-t 1` goes from exit 1 to exit 0, and the full
default sweep now completes with exit 0. Max err_aiter per table --
biased_grouped_topk 1.5e-05, kimi_k25 1.3e-05, kimi_k3_strided 3.7e-04,
grouped_topk 0.125 -- all as before, no AssertionError anywhere.

Not a regression: this branch never touched the block (git diff against
origin/main reports no change to it), and `-t 1` fails identically on
origin/main. It arrived with 3135022 (#4397), the same PR as the mxfp4
scale-pad reference bug fixed in the previous commit.

Swept the other `assert not X.is_contiguous()` sites -- 9 of them, in
test_fused_qk_norm_rope_cache_quant.py, triton_tests/attention/
test_mha.py and test_mha_v3.py. None of them slices a dimension that can
collapse to 1 the way this one does.
valarLip added a commit that referenced this pull request Jul 28, 2026
* [CI] pin ruff and make its configuration explicit

The style job runs `pip3 install ruff` unpinned against a repo with no ruff
config, so the enforced rule set is whatever the latest release happens to
default to. ruff 0.16.0 widened those defaults well past E4/E7/E9/F, which
turned PRs red for reasons that had nothing to do with their contents. This
pins the inputs; it does not change any verdict on existing code.

- Pin ruff==0.16.0 in .github/workflows/pre-checks.yaml. Bumping it is now the
  single reviewable place where the rule set can move.
- Add [tool.ruff] to pyproject.toml with target-version = "py310". There is no
  [project] table, so ruff cannot read requires-python and was falling back to
  its own default; py310 is what it already resolved to, so the UP0xx verdicts
  are unchanged. Deliberately no select/ignore override -- the rule set stays
  exactly ruff's default for the pinned version.
- Restate the default exclude list with `dist/` instead of the bare `dist`. A
  pattern without a slash matches the basename at any depth, so the built-in
  entry was silently skipping aiter/dist/ -- 15 real source modules that had
  never been linted. They contribute 346 pre-existing findings, to be burned
  down with the rest of the backlog. The setuptools output directory stays
  excluded. Also exclude 3rdparty/, aiter/jit/build/ and aiter_logs/.
- setup.py: python_requires ">=3.8" -> ">=3.10". 3.8/3.9 have not worked for a
  long time -- 81 modules use PEP 604 annotations without
  `from __future__ import annotations` and raise TypeError at import on <3.10.
- bench_deepgemm_attention.py: the AOT name f-string reused its outer quotes,
  which is PEP 701 (3.12+) and left the module unparseable on 3.10/3.11. Use
  single quotes inside. Found by pinning target-version.

Verified against main: no existing finding disappears (nothing is weakened) and
no new finding appears outside aiter/dist/. black --check stays clean.

* fix: accept PEP 604 unions in generate_schema

generate_schema matched Tensor-bearing unions with `get_origin(t) is Union`,
but `X | None` has origin types.UnionType, not typing.Union. Every PEP 604
annotation therefore fell through to the "unsupported type" branch and the
parameter was silently dropped from the registered custom-op schema:

    def f(a: torch.Tensor, b: Union[torch.Tensor, int] = None)
      -> (Tensor(a0!) a, Tensor(a1!)? b=None) -> Tensor
    def f(a: torch.Tensor, b: torch.Tensor | int = None)
      -> (Tensor(a0!) a, * ) -> Tensor          # b is gone

This is pre-existing: 449 annotations across 99 modules already use PEP 604.
aiter/jit/core.py got this right via its _is_union() helper; torch_guard.py is
the one place that did not. Match core.py's behaviour.

* style: apply ruff's mechanical autofixes (batch 1/N)

Safe autofixes only, for rules where every finding in the repo was
safe-fixable. No behaviour change:

  RUF100   305  unused noqa directives
  PIE808    87  range(0, n) -> range(n)
  PLR0402   65  import x.y as y -> from x import y
  PLR2044   59  empty comments
  PIE790    46  unnecessary pass / ...
  RUF010    42  f"{str(x)}" -> f"{x!s}"
  C408      51  dict() -> {}   (ruff marks this fix unsafe; every call site
                 uses identifier keys and `dict` is not shadowed anywhere in
                 the repo, both checked)

Applied with `--fixable`, not `--select`. That distinction matters: `--select`
narrows the *enabled* rule set, so RUF100 then treats every noqa referencing a
non-selected rule as unused and strips it. A first attempt that way deleted
live suppressions -- `# noqa: F821` on gfx950_a8w8_blockscale_asm (a name the
compile_ops decorator injects at runtime) and several `# noqa: E722` -- and
introduced 51 new findings. `--fixable` keeps the full rule set enabled for
analysis and only applies the listed fixes.

Verified: ruff total 6678 -> 6022, and the per-rule diff is reductions only --
no rule gains a single finding. black --check clean, all 261 files compile.

* style: modernize type annotations (batch 2/N)

PEP 585 / PEP 604 annotations, now that python_requires is >=3.10:

  UP045  2165  Optional[X] -> X | None
  UP006   613  typing.Dict/List/Tuple -> dict/list/tuple
  UP007   109  Union[X, Y] -> X | Y
  UP035   166  Callable/Iterable/Iterator -> collections.abc
  F401    410  typing imports left unused by the above

Applied with `--fixable` so the full rule set stays enabled during analysis
(see batch 1 for why `--select` is wrong here). None of the unused imports
were in an __init__.py, so no re-export was dropped.

Two real bugs this surfaced, both fixed here:

- Optional[torch.tensor] -- lowercase `tensor`, the factory function, not the
  `Tensor` type -- in fused_moe.py, fused_moe_dp_shared_expert.py,
  test_common.py and test_pa_prefill.py. typing.Optional accepts any callable
  and silently swallowed it; `torch.tensor | None` raises TypeError at import.
  Corrected to torch.Tensor. Without this, `import aiter.fused_moe` fails
  outright, which is how it was caught.
- The generate_schema union handling, fixed in the preceding commit; it had to
  land first or every converted Tensor union would have produced a truncated
  custom-op schema.

Verified:
- Import sweep over every aiter submodule, before vs after: 249 ok / 1 fail
  both times, the one failure (pod_attention) pre-existing and unrelated.
- generate_schema returns identical schemas for Optional[T], T|None,
  Union[T,int] and T|int.
- ruff total 6022 -> 2936, per-rule diff is reductions only.
- black --check clean, all 252 files compile.

* style: finish the batch 2 residue that ruff could not autofix

Ten findings ruff left behind, now resolved so UP006/UP007/UP035/UP045 and the
batch 1 rules are all at zero:

- dist/utils.py: the recursive JSONTree alias -> PEP 604. Verified the runtime
  form actually builds (forward-ref strings inside `|` are fine).
- rotary_embedding.py: List[float] -> list[float] and drop the typing.List
  import. The bare `# noqa` on that line was suppressing UP006 and is no longer
  needed -- bare noqa hides everything, so removing it restores coverage.
- torch_guard.py: the three `param_type == Optional[X]` comparisons are runtime
  values rather than annotations; both spellings compare equal, so converting
  them is safe and the comment now says so.
- torch_guard.py / triton/__init__.py: `# noqa: F401` with a reason on the
  `Library` availability probe and on the deliberate comms re-exports.
- Two typing imports left unused by the above.

Note for reviewers: check residues with the full rule set, never with
`--select`. `--select` narrows the *enabled* rules, so RUF100 then reports every
noqa referencing an unselected rule as unused -- it flagged 29 phantom findings
here, the same trap that corrupted the first attempt at batch 1.

Re-verified after the change: import sweep still 249 ok / 1 pre-existing fail,
generate_schema identical across spellings, black clean, all files compile.

* style: sort imports where ordering is not load-bearing (batch 3a/N)

I001 across the 620 files whose module-level imports all sit at the top of the
file with no sys.path manipulation. The other 43 I001 files -- code generators
under csrc/ that insert into sys.path before importing, tests doing the same,
and the runtime's deliberate late imports -- are left for a follow-up that has
to read each one.

Static "imports are all at the top" is necessary but not sufficient: it says
nothing about whether their relative order matters. aiter/__init__.py proved
that. Sorting its block moved `from . import mla` ahead of
`from .utility import dtypes`, and mla.py does `from aiter import dtypes`, so
`import aiter` died with a circular-import ImportError. That block is now
fenced with `# isort: off` / `# isort: on` and a comment explaining why the
order is load-bearing; the file's other two import blocks are sorted normally.

Caught by importing every aiter submodule and diffing against the pre-change
baseline -- 249 ok / 1 pre-existing failure both before and after, with the
regression showing up as a hard failure of `import aiter` itself. ruff total
2926 -> 2262 (I001 754 -> 90), no other rule moves in either direction.
black --check clean, all 620 files compile.

Not covered: op_tests/ and csrc/ scripts are not exercised by the import sweep,
so their import-order behaviour rests on CI.

* style: clear the pre-existing findings this PR pulled into the diff

CI's reviewdog runs with filter-mode=diff_context, so touching 756 files
dragged ~117 findings that were already in the tree into review scope. None of
them are introduced by this PR; they are simply adjacent to lines it changed.
Clearing them is what it takes to get the PR green.

Autofixed: PYI041, RUF022, RUF013, RUF059, SIM114, SIM118, SIM210, C417,
FURB136, PLR1711, PLC0206, PLW1508, LOG015, I001.

Real fixes:
- cuda_wrapper.py: two mutable class attributes -> ClassVar (RUF012).
- quick_all_reduce.py: dropped `from aiter import logger`, immediately shadowed
  by the module's own getLogger (F811).
- rotary_embedding.py: os.environ.get(..., 0) -> "0" -- the int default was
  never reachable but is the wrong type (PLW1508).
- base_tuner.py, test_pa_decode_gluon.py, bench_fav3_sage.py: unused unpacked
  values, a mutable default argument, and a dead self-assignment.
- flydsl_dispatch_combine: bind expert_rsrcs/expert_vlds as closure defaults
  (B023), same pattern as the other FlyDSL kernels.
- 13 files: cleared a stray executable bit; one file got the bit it needed.
- 4 op_tests modules keep their camelCase names with a file-level noqa; renaming
  them would break the test-selection lists that reference those paths.

`# noqa` with a stated reason where the finding is a deliberate choice:
intentional blanket catches, framework defaults evaluated at import,
PEP 484 redundant-alias re-exports, and a best-effort probe.

Third instance of the same latent bug: RUF013 rewrote `x: callable = None` into
`callable | None`, and `callable` is the builtin function, not a type -- so
`import aiter` died again, exactly like Optional[torch.tensor] did in batch 2.
Fixed in all 6 files to collections.abc.Callable. The annotation was always
wrong; making unions real is what turns it into a hard failure.

Verified: diff-context findings 117 -> 0 (measured locally against merge-base,
since `gh pr diff` truncates a PR this size), import sweep back to the 249 ok /
1 pre-existing fail baseline, black clean, all touched files compile.

* style: apply every remaining ruff autofix (safe and unsafe)

2076 -> ~830 findings. Everything ruff can rewrite mechanically, plus the
executable-bit and shebang findings:

- EXE001/EXE002: 34 files got the executable bit their shebang asks for, 15
  lost a bit they had no shebang for.
- EXE005: pa_mqa_logits.py's 66 `#!=----` separators are decorative comments,
  not shebangs -- file-level noqa rather than rewriting them all.
  hipify_python.py had a real shebang below the licence header; moved to line 1.
- RUF013's 38 sites were pre-checked for the trap that bit batch 2 twice: every
  annotation involved is a genuine builtin type (int/float/str), no factory
  functions, so making the Optionals explicit is safe.

One autofix had to be reverted: TC004 hoisted `from fla import __version__`
out of its `if TYPE_CHECKING:` block in gated_delta_rule_utils.py. fla is an
optional dependency, so that turned `import aiter.ops.triton.gated_delta_net`
into a ModuleNotFoundError. The guard is restored with a noqa explaining why.
(The runtime read of __version__ under a TYPE_CHECKING-only import is a
pre-existing latent bug in this vendored file; not touched here.)

Import sweep back at the 249 ok / 1 pre-existing fail baseline, black clean,
all 357 touched files compile.

* style: mechanical fixes for the rules ruff cannot autofix

- LOG015 (145): five files logged through the root logger; each now has a
  module-level getLogger(__name__) and uses it. Only the logging calls were
  rewritten, not basicConfig/level constants.
- RUF012 (30): mutable class attributes annotated ClassVar.
- PLW1508 (57 total): os.environ.get/getenv integer defaults quoted; the value
  was always passed through int() so behaviour is unchanged.
- RUF059 (27): unused unpacked values prefixed with an underscore.
- PLW0602 (20): dropped 'global' statements that only ever read the name.
- PLW1510 (17): explicit check=False on subprocess.run -- this is the existing
  default, so no behaviour change.

Import sweep still 249 ok / 1 pre-existing fail; black clean; everything
compiles.

* fix: correct two misplaced noqa directives

- gated_delta_rule_utils.py: TC004 is reported on the import line, not on the
  `if TYPE_CHECKING:` header, so the noqa sat one line too high and RUF100
  flagged it as unused while TC004 still fired.
- bench_attn_models.py: the LOG015 noqa became dead once this file was switched
  to a module logger in the LOG015 pass.

Both were self-inflicted in earlier commits on this branch.

* style: mark the deliberate blanket catches with noqa (BLE001)

All 186 BLE001 sites across 102 files. Sampling them shows one consistent
shape rather than sloppiness, so suppressing beats rewriting:

- tuners and benchmark sweeps that must record a failed shape and keep going
  instead of aborting the whole run
- optional fast-path probes (fused kernels, ipex, iris, fla) that fall back to
  a reference path when the import or the call fails
- best-effort cleanup and environment introspection

Narrowing these to specific exception types would need to enumerate everything
a GPU kernel launch, a JIT build or an optional dependency can raise, which is
exactly the open set the blanket catch exists to absorb.

Deliberately a bare `# noqa: BLE001` per site rather than a repo-level ignore:
the rule keeps applying to new code, and a reviewer sees the suppression at the
point of use. The rationale lives here rather than repeated 186 times.

One site needed care: black had reflowed `except Exception as e:` in
tuned_gemm.py into a parenthesised multi-line form once the comment pushed it
past the line limit, which stranded the noqa on the wrong line. Rewritten as a
single-line except with the explanation moved above it.

Import sweep 249 ok / 1 pre-existing fail, black clean, 102 files compile.
Repo total 566 -> 380; the remaining rules are still open for review.

* style: bind loop variables in closures (B023)

All 106 sites across 30 files, 42 distinct closures (93 in def, 12 in lambda),
fixed properly rather than suppressed: the loop variable is now bound as a
default argument, so each closure captures its own iteration's value.

    def _accum_loop(end, U):            ->  def _accum_loop(end, U, hdim_off=hdim_off):
    lambda: bench(x_aiter)              ->  lambda x_aiter=x_aiter: bench(x_aiter)

This is the right fix here rather than a noqa. Most of these are FlyDSL/Triton
codegen closures that are invoked inside the same iteration, so late binding
happens to give the right answer today -- but nothing enforces that, and the
binding makes the intent explicit. Where a closure genuinely outlives its
iteration, B023 was flagging a real bug and this fixes it.

Verified as signature-only: stripping the newly added `name=name` defaults from
both sides makes the before and after ASTs compare equal across all 30 files,
so no function body changed. The closures have no *args/**kwargs/kw-only
parameters, so appending defaulted parameters cannot shift an existing binding.

Two bugs in my own transformation, caught before anything was written:
- zero-argument lambdas are spelled `lambda:` with no space, so inserting the
  binding before the colon produced `lambdax=x:`. The script re-parses every
  file it edits and skipped the three affected files instead of corrupting them.
- ruff reports B023 at the *use* site, so the enclosing closure has to be
  resolved by innermost-enclosing-node rather than by line proximity.

Import sweep 249 ok / 1 pre-existing fail, black clean, 30 files compile.
Repo total 380 -> 274.

* style: make the warmup group access explicit (B018)

Both sites are the same line in a multi-GPU test:

    # warmup and align all gpu
    get_tp_group().device_group

device_group is a plain attribute assigned in GroupCoordinator.__init__, not a
lazy property, so the attribute access really is a no-op -- B018 is right about
that half. What the statement still does is call get_tp_group(), which raises if
the TP group was never initialised, so it is not dead code either.

Bound to `_` and documented rather than deleted, which keeps that check and
satisfies the rule. Both files are otherwise ruff-clean now.

* style: clear S110/B008/TRY002/SIM102/PLR0124/PLW0127/TRY004

Real fixes:
- TRY002 (33): `raise Exception(...)` -> `raise RuntimeError(...)`. RuntimeError
  is an Exception subclass, so every existing `except Exception` still catches.
- SIM102 (38): nested ifs merged with `and`, which is exactly equivalent when
  neither branch has an else. 35 done by an AST-driven pass, two more by hand
  (a three-level nest, and one whose inner condition carried an inline comment).

noqa where the finding is deliberate:
- S110 (30): try/except/pass around optional imports and best-effort probes,
  already carrying `# noqa: BLE001` -- S110 merged into the same directive.
- B008 (41): framework defaults such as fx.Stream(None) and Range(...) that are
  meant to be built once at import.
- PLR0124 (21): `x == x` is the NaN test in Triton kernels. Rewriting it as ruff
  suggests would delete the check.
- PLW0127 (13): `accumulator = accumulator` keeps a Triton if/elif chain
  structurally uniform.
- TRY004 (12): ValueError is the established contract of these helpers.

Three SIM102 sites were left nested on purpose, each with the reason inline. One
matters: `# pyre-ignore[61]` sits between the two ifs in hstu_attention.py and
applies to the inner one, so merging would have silently dropped a type-checker
suppression.

Placing noqa needed a convergence loop rather than a single pass: adding one can
push a line past the limit, black then splits the statement, and the directive
ends up on a line where the rule no longer fires -- RUF100 flags it while the
original finding comes back. The loop adds, reformats, strips whatever RUF100
calls unused, and repeats.

Import sweep 249 ok / 1 pre-existing fail; black clean; 112 files compile.
Repo total 272 -> 84.

* style: drop the fx.Stream(None) parameter defaults (B008)

    stream: fx.Stream = fx.Stream(None)  ->  stream: fx.Stream

All 33 FlyDSL jit launchers. Replaces the noqa these carried a moment ago: the
default really was being built once at import and shared, which is what B008
objects to.

Checked every call site before removing it, since this makes the parameter
required. 17 pass a stream (the `_run_compiled(fn, (..., torch.cuda
.current_stream()))` form), and the four that looked like they did not:

- two `launch()` calls in aot/flydsl/common.py are a different, local
  scheduling helper -- name collision, not these functions
- flydsl_fp8_mqa_logits takes a plain `stream=None`, a different pattern that
  was never part of this set
- launch_pa_decode_ps_reduce_flydsl in pa_decode_gluon.py genuinely relied on
  the default. That one call site now passes `stream=fx.Stream(None)`
  explicitly, which is byte-for-byte the previous behaviour (fx.Stream(None) is
  the default queue) and moves the construction out of the signature.

`stream` is the last parameter in all 33, so removing its default cannot produce
a non-default-after-default signature; verified via AST rather than by eye.

Import sweep 249 ok / 1 pre-existing fail, black clean, 40 files compile.

Caveat: the import sweep proves these modules still load, not that a kernel
still launches. The launcher call paths need GPU CI.

* style: clear F821, SIM115 and UP031

F821 (12) -- three different causes, so three different answers:
- IrisCommContext was referenced from a string annotation in all_gather.py and
  reduce_scatter.py but never imported. It does exist, in comms/iris.py, so
  both files now import it under `if TYPE_CHECKING:`.
- test_pa_ragged_experimental.py is a near-copy of csrc/cpp_itfs/pa/
  pa_ragged_test.py that lost STR_DTYPE_TO_TORCH_DTYPE along the way. Its
  get_kv_cache_torch_dtype would NameError on any string cache_dtype. Dict
  restored from the original.
- The rest are unfixable here and got a noqa plus a file-level note: dist/
  utils.py is vendored from vLLM and reaches for `envs`, `current_platform` and
  `enable_trace_function_call` from vllm modules that were never carried over,
  and mla_extend_ref.py calls vLLM's context_attention_fwd, which was never
  ported. Adding an import is not an option; these paths simply cannot run.

SIM115 (8):
- Three real leaks-by-omission converted to `with`: view-screen.py held a json
  handle across a 50-line loop, test_common.py opened a metafile inside a list
  comprehension, test_gemm_codegen.py read a file without closing it.
- Five are `tempfile.NamedTemporaryFile(delete=False)`, where the path
  deliberately outlives the handle. noqa'd with that reason.

UP031 (8): printf-style formatting to f-strings. Seven are the same
"%d / %d mismatched elements" line duplicated across the moe tests.

Import sweep 249 ok / 1 pre-existing fail, black clean, 18 files compile.
Repo total 55 -> 27.

* style: clear the last 27 findings -- repo is ruff-clean

Four of these were real bugs, not style:

- PLE1205: logger.warning("Custom quick allreduce:", f"Invalid quantization
  level: ...") passed the detail as a %-arg to a format string with no
  placeholders, so it never reached the log. Joined into one message.
- F601: a8w8_gemm1_kernels_list_gfx950 keyed two entries `4`, so the first was
  silently dropped at dict construction and key 3 did not exist. Renumbered the
  first to 3. Lookups for 1/2/4/5/6 are unaffected -- 4 already resolved to the
  second entry -- so this only restores the missing 3. Worth an author check
  that 3 is indeed the intended key.
- RUF034: two `x.stride(0) if not flash_layout else x.stride(0)` in
  fused_kv_cache.py, identical on both branches while the neighbouring lines
  genuinely permute. Collapsed; flagging in case the layout swap was meant to
  differ here too.
- B012: three `return` statements inside `finally` in test_communication.py,
  which swallow anything the handler above re-raises. Moved out of the finally.

Behaviour-preserving cleanups: dead `run_aiter_hip` in test_pa_mtp.py (shadowed
by a later definition, so the removed one never ran) and its now-unused import;
the unused `remap_xcd` import in lean_atten.py that a local @triton.jit
definition overrode; SIM103/SIM113/PERF402/PLC0206/E722/DTZ005/UP031 rewrites.

Two renames needed scope checks before I trusted them. get_args_of_build's loop
rebinds its own `ops_name` parameter, and `ops_name` is read again at three
later lines -- but those sit in elif/else branches that the loop's branch
returns before reaching, so the shadowing is unreachable and the rename is
cosmetic. test_la_paged's `for h in range(h)` shadows the head-count argument;
every `h` in the body is therefore the index, so renaming the loop and its body
to `hi` is exact.

noqa with a reason for the two that should stay: functools.lru_cache on a
method of a process-lifetime singleton (B019), and the exec() the JIT loader
needs (S102).

`ruff check .` now passes on the whole repository. black clean, import sweep
249 ok / 1 pre-existing fail, all touched files compile.

* fix: restore the torch import that F401 removed from core.py's eval path

CI caught this: `python setup.py` died with

    File "aiter/jit/core.py", line 1145, in convert
      val[idx] = eval(el)
    NameError: name 'torch' is not defined

get_args_of_build eval()s strings out of optCompilerConfig.json, some of which
reference torch. The original code bound the name right before the eval:

    if "torch" in el:
        import torch as torch
    val[idx] = eval(el)

The `as torch` spelling was almost certainly there to mark the import as
deliberate. F401 does not read eval() strings, saw an unused import, deleted it,
and left `pass` behind -- so every build that hit a torch-bearing config entry
crashed. Restored as a plain `import torch` with a noqa and a comment saying
what the binding is for.

Swept for the same failure mode across the branch: compared the set of imported
names before and after per file, kept only names that disappeared entirely (a
line that merely lost a `# noqa` is not a removal), intersected with files that
actually call eval/exec, and checked whether the name still appears in a string
literal. One candidate came back, `List` in core.py, and it is a false positive
-- the exec() namespace at line 1630 maps "List" explicitly, it does not read
the module scope. This was the only real one.

Verified by reproducing the CI command: `python setup.py --help` now completes,
and get_args_of_build("all") returns its 106 srcs. ruff and black still clean.

* fix: make ck_tile_gemm_moe_2stages codegen set up sys.path before importing

gen_instances.py had all its imports at the top of the file and relied on
an import-time side effect of a sibling: moe_cktile2stages_common does a
sys.path.insert() for aiter/jit/utils, and that is what made `chip_info`
resolvable. isort sorted the two local imports alphabetically, putting
chip_info first, so the blob_gen_cmd subprocess died with:

    File "csrc/ck_tile_gemm_moe_2stages/gen_instances.py", line 11
    from chip_info import get_gfx, get_gfx_list
    ModuleNotFoundError: No module named 'chip_info'

The file was the odd one out. Every other chip_info consumer in csrc/
(10 of them) puts the sys.path.insert() above its local imports, which
both makes them self-sufficient and gives isort a barrier it will not
sort across. Give this file the same shape instead of fencing it with
`# isort: off`: the hidden cross-file coupling goes away, the import
block is free to be sorted, and the file now matches its siblings.
Its own sys.path.insert() -- previously sitting below the imports, where
it did nothing for them -- simply moves up.

Verified: the real blob_gen_cmd (`gen_instances.py --working_path ...`)
exits 0 and emits 224 files, byte-identical to the pre-regression output.

Swept for the same class repo-wide:
  - top-level import order before/after across all 821 changed .py files
    vs every module that mutates sys.path at import time (20): 1 hit,
    this file.
  - all 17 blob_gen_cmd codegen scripts from optCompilerConfig.json run
    end to end: all pass (hsa/codegen.py for each of -m pa/mla/mla_v4/
    i8gemm/bf16gemm, once AITER_GPU_ARCHS is set).
  - the 300 reordered scripts outside aiter/, differentially probed
    old-vs-new: 264 clean, 36 failures that reproduce identically on the
    base revision, 0 regressions.
  - repo-wide scan for the fragile shape itself (an injector imported
    inside an unbarriered leading import block): 2 files, this one and
    ck_gemm_moe_2stages_codegen/gen_instances.py, whose only other
    imports are stdlib and is therefore unaffected.

* fix: bind Tuple explicitly in the check_args exec namespace

check_args() rewrites a pybind op's __doc__ into a `def` and execs it to
recover the signature. The rewrite emits List[, Tuple[ and Optional[, but
only List and Optional were bound in the exec namespace -- Tuple happened
to be supplied by the `from aiter import *` on the first line of the
exec'd source, because some submodule under aiter did `from typing import
Tuple` at module scope and star-import leaked it into the aiter namespace.

UP035 rewrote every `from typing import Tuple` in the repo, so the leak
stopped and the first call to any op with a tuple in its signature died:

    File "aiter/jit/core.py", line 1637, in check_args
      exec(  # noqa: S102
    File "<string>", line 2, in <module>
    NameError: name 'Tuple' is not defined

Bind "Tuple": tuple next to "List": list. Star-import was never a sound
source for these; the comment now says so.

Verified: `Tuple in dir(aiter)` is True on the base revision and False at
HEAD; replaying the exec with the old namespace reproduces the CI
NameError verbatim and with the new one it resolves.

Swept the whole class rather than this one name: dumped the star-import
surface of all 322 importable aiter modules on both revisions and diffed
them (aiter_export_surface_diff.py).

  - top-level `aiter` lost exactly 3 names -- List, Optional, Tuple. The
    first two were already bound explicitly, so Tuple was the only gap.
    Nothing was gained, and no module became unimportable.
  - across all modules, 210 names lost, 203 of them typing re-exports.
    The other 7 (logging, get_hip_arch, torch, lru_cache, tl, triton)
    belong to modules nothing star-imports; aiter/__init__.py is the only
    star-importer in the repo and its result is the surface already
    diffed above.
  - enumerated every eval/exec/compile site in the repo and intersected
    with the per-module lost names: only core.py (lost List) and
    tuned_gemm.py (lost Optional) qualify. core.py's evals consume
    optCompilerConfig.json, where "List" appears 0 times, and its exec
    binds List explicitly; tuned_gemm.py's evals consume dtype strings
    from configs/*.csv, where "Optional" appears 0 times. Both clear.

* fix: restore the paged_attn binding test_pa_mtp lost to F401

Standard Tests shard 1 died with:

    File "op_tests/test_pa_mtp.py", line 272, in run_aiter_hip
      return aiter.paged_attn.PagedAttention.forward_decode(
    AttributeError: module 'aiter' has no attribute 'paged_attn'

Two rules interacted. The file used to define run_aiter_hip twice; the
second shadowed the first, so the first was dead code and F811 removed
it. The dead copy was the only user of the `ops` alias from
`from aiter import dtypes, paged_attn as ops`, so F401 then removed that
import as unused -- correct in isolation, except the import was also
what bound `paged_attn` as an attribute of the aiter package, which is
how the surviving copy reached it. Neither rule can see that side effect.

Import paged_attn under its own name and call it directly instead of
going through the package attribute, so the dependency is visible to
static analysis and needs no noqa.

Verified: `aiter.paged_attn` does not resolve after a plain `import
aiter`, which is exactly why the attribute access failed; running
op_tests/test_pa_mtp.py now executes into aiter/paged_attn.py:272 and
only stops later at a JIT compile step, because git worktree does not
populate 3rdparty/composable_kernel.

Swept both rules across the branch:
  - every import removed from the 821 changed files, checked for one
    whose parent package is still reached by attribute access: this file
    was the only true hit. test_sample.py looked similar but
    aiter.greedy_sample / aiter.random_sample resolve on a plain import
    anyway.
  - widened to every `aiter.X` access in the changed files, resolved
    against a real import: 11 do not resolve, and all 11 fail the same
    way on the base revision, so none of them are new.
  - definitions removed anywhere in the branch: exactly one, the
    shadowed run_aiter_hip above.

The other failure in that shard, test_quant_mxfp4.py
test_gui_shuffle_scale_w2_k_pad inter_tp=640, is not ours: the same job
on main at ad2f051 fails on the same assertion at the same line.

* fix(test): pad the mxfp4 GUI scale reference with the neutral e8m0 scale

test_gui_shuffle_scale_w2_k_pad has been failing on main at inter_tp=640:

    File "op_tests/test_quant_mxfp4.py", line 334
      assert torch.equal(
    AssertionError: inter_tp=640

The kernel is fine; the reference model in the test is wrong. Both sides
of the assert call the same shuffle_scale -- the only difference is who
pads k_groups up to a multiple of 8. shuffle_scale does it internally
with 0x7F for 1-byte scales (aiter/ops/shuffle.py:372), which in e8m0 is
the exponent bias, i.e. 2^0. The test pre-padded with torch.zeros, and
0x00 in e8m0 is 2^-127, so the two layouts genuinely differ.

inter_tp=256 gives k_groups=8, already a multiple of 8, so no padding
happens and the fill value never matters -- which is why only the 640
case failed.

Pad the reference with the neutral scale instead, mirroring what the
implementation guarantees: k-groups added by padding describe weights
that do not exist, so their scale has to be 1.0.

Verified on gfx950: with torch.zeros the assert fails for both float16
and bfloat16 at inter_tp=640 and passes at 256; with the neutral fill
all four combinations are byte-identical. `python3
op_tests/test_quant_mxfp4.py --a16w4_shuffle` goes from exit 1 to exit
0 on the same tree, and the full file passes.

Not a regression -- the test was written this way. The internal padding
and its 0x7F fill both landed in c2b69b9 (#4130, 2026-07-09);
test_gui_shuffle_scale_w2_k_pad was added later in 3135022 (#4397,
2026-07-28) and did not mirror the fill value.

Swept for the same assumption elsewhere: no other test hand-pads a scale
tensor before calling shuffle_scale. Every other caller passes unpadded
production data and relies on the internal padding, which is correct.

* fix(test): drop the non-contiguity assert in the strided grouped-topk sweep

Running op_tests/test_moeTopkSoftmax.py with its default token list dies
immediately:

    File "op_tests/test_moeTopkSoftmax.py", line 708, in <module>
      assert not gating_output.is_contiguous()
    AssertionError

The block builds a row-strided view of a fused router buffer:

    backing = torch.randn((token, 6016), dtype=bf16)
    gating_output = backing[:, 1536:1536 + 896]

and the default token list starts at 1. At token=1 the view is
(1, 896) with stride (6016, 1), and a row stride you can never take is
not a stride at all -- torch reports it as contiguous. The assert is
false for the first case it ever sees.

`assert gating_output.stride(0) == fused_front_width` is the check that
actually pins the layout under test, and it holds for every token. Keep
it and drop the other one.

Verified on gfx950: `-t 1` goes from exit 1 to exit 0, and the full
default sweep now completes with exit 0. Max err_aiter per table --
biased_grouped_topk 1.5e-05, kimi_k25 1.3e-05, kimi_k3_strided 3.7e-04,
grouped_topk 0.125 -- all as before, no AssertionError anywhere.

Not a regression: this branch never touched the block (git diff against
origin/main reports no change to it), and `-t 1` fails identically on
origin/main. It arrived with 3135022 (#4397), the same PR as the mxfp4
scale-pad reference bug fixed in the previous commit.

Swept the other `assert not X.is_contiguous()` sites -- 9 of them, in
test_fused_qk_norm_rope_cache_quant.py, triton_tests/attention/
test_mha.py and test_mha_v3.py. None of them slices a dimension that can
collapse to 1 the way this one does.
apicciau added a commit that referenced this pull request Jul 28, 2026
* Resolve conflicts with Kimi-K3 SiTUv2 MoE PR (#4397) and ruff formatting (#4403)
* Keep both a16w4 INTERLEAVE-gate additions and upstream SiTUv2/ruff changes
* Deduplicate imports in moe_gemm_2stage.py
* Coexist include_stage2_out and pass_swiglu_limit in _s1_args_fp4
@zufayu
zufayu deleted the k3-for-amd branch August 11, 2026 06:04
zufayu added a commit that referenced this pull request Sep 7, 2026
FlyDSL needs patching after merge 1.66x as often as the repo average: 20 of
the 98 fix-commits that reference another PR touch it, against a 12.3% share
of all commits (194/1577), one-sided binomial p=0.015, measured on aiter main
636098e. Its revert rate is NOT elevated -- 5 of 40, against that same
12.3% -- but n=40 has no power in either direction, so that is "cannot tell",
not "no". The patching is where the signal is, and post-merge patching is
what review is for.

Of its 21 precise fixes, 9 fall into two mechanisms and neither had a rule.

B8 -- a buffer or descriptor bound not tied to the tensor's real extent.
de9f1f8 bound arg_scale_a with max_size=True while scale_a is per-row(M)
and M is ragged; the comment it also corrected ("B and scales are
exact-multiple in N") was false for scale_a, so the code was written to match
a premise that did not hold. Nothing covered this class: num_records,
max_size, buffer_resource and oob_ appear nowhere in rules.md, and B2 -- the
one out-of-bounds rule -- is tl.load without a mask, which cannot fire on a
backend that has no tl. A FlyDSL PR deriving memory-safety got B2 D1 G1.

D12 -- a new ops-side contract the AOT pre-compiler was not taught. #4397
added resolve_flydsl_stage1_tile_n and resolve_flydsl_stage2_tile_k in one
commit, wired only stage2 into AOT, and the stage1 drift lived 32 commits
until #4429.

Both collectors were checked against the commits that introduced the defects,
not only against fixtures. flydslbounds reports arg_scale_a on 3910ab3 and
falls silent on de9f1f8, so the fix does not trip its own rule. aotpair
reports resolve_flydsl_stage1_tile_n on 3135022 and stays quiet about
stage2 -- which that same commit did wire in -- which is the distinction the
rule exists to make.

The first cut of D12 looked for "ops changed, AOT untouched" and would have
found nothing: #4397 does touch AOT (grouped_moe.py +19, moe.py +9). The
defect is a path missed, not a side skipped.

Separately, the card gate had exactly one way through and it was dishonest.
On wave5 PR4862 an agent found a CI problem at Step 7 that no rule covered,
the gate called it UNBACKED, and the only move that clears that gate is for
the finding to have been known earlier -- so it edited answers.txt, a file
that had already passed its own gate, and described this itself as "slightly
dishonest bookkeeping". A gate whose only path through is fabrication trains
fabrication. `-- late finding:` backed by an append-only late_findings.txt is
the exit; a hash ledger keyed on the FIRST passing hash is the lock. Neither
half works alone: the exit alone leaves the door open, the lock alone leaves
a reviewer who genuinely found something late with nowhere to put it.

Two changes planned from reading the code and dropped after measuring them:
adding .cpp to is_kernel (aiter has 12 .cpp/.cc files and none contains a
device kernel, so excluding them is correct) and dropping the .py test in
flydsl-kernel (non-Python is deliberately left to the C path, and a test
says so).
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

10 participants