Skip to content
Merged
Show file tree
Hide file tree
Changes from 6 commits
Commits
Show all changes
25 commits
Select commit Hold shift + click to select a range
0a81a5c
[CI] pin ruff and make its configuration explicit
valarLip Jul 28, 2026
ad8db44
fix: accept PEP 604 unions in generate_schema
valarLip Jul 28, 2026
e40c1cb
style: apply ruff's mechanical autofixes (batch 1/N)
valarLip Jul 28, 2026
b23a643
style: modernize type annotations (batch 2/N)
valarLip Jul 28, 2026
ed8e227
style: finish the batch 2 residue that ruff could not autofix
valarLip Jul 28, 2026
464dd84
style: sort imports where ordering is not load-bearing (batch 3a/N)
valarLip Jul 28, 2026
a459af7
style: clear the pre-existing findings this PR pulled into the diff
valarLip Jul 28, 2026
da4ba79
style: apply every remaining ruff autofix (safe and unsafe)
valarLip Jul 28, 2026
9b0975b
style: mechanical fixes for the rules ruff cannot autofix
valarLip Jul 28, 2026
ffde60b
Merge origin/main into ruff/pin-and-configure
valarLip Jul 28, 2026
6f02284
fix: correct two misplaced noqa directives
valarLip Jul 28, 2026
b652709
style: mark the deliberate blanket catches with noqa (BLE001)
valarLip Jul 28, 2026
a15e9fb
style: bind loop variables in closures (B023)
valarLip Jul 28, 2026
64c4228
style: make the warmup group access explicit (B018)
valarLip Jul 28, 2026
2f3ce97
style: clear S110/B008/TRY002/SIM102/PLR0124/PLW0127/TRY004
valarLip Jul 28, 2026
266c97f
style: drop the fx.Stream(None) parameter defaults (B008)
valarLip Jul 28, 2026
596c9d0
style: clear F821, SIM115 and UP031
valarLip Jul 28, 2026
c9af09e
style: clear the last 27 findings -- repo is ruff-clean
valarLip Jul 28, 2026
4a68321
fix: restore the torch import that F401 removed from core.py's eval path
valarLip Jul 28, 2026
0d763d2
fix: make ck_tile_gemm_moe_2stages codegen set up sys.path before imp…
valarLip Jul 28, 2026
7d4518c
fix: bind Tuple explicitly in the check_args exec namespace
valarLip Jul 28, 2026
0a4acd2
Merge branch 'main' into ruff/pin-and-configure
valarLip Jul 28, 2026
e688cb5
fix: restore the paged_attn binding test_pa_mtp lost to F401
valarLip Jul 28, 2026
783dc22
fix(test): pad the mxfp4 GUI scale reference with the neutral e8m0 scale
valarLip Jul 28, 2026
092427c
fix(test): drop the non-contiguity assert in the strided grouped-topk…
valarLip Jul 28, 2026
File filter

Filter by extension

Filter by extension


Conversations
Failed to load comments.
Loading
Jump to
The table of contents is too big for display.
Diff view
Diff view
  •  
  •  
  •  
19 changes: 9 additions & 10 deletions .github/scripts/compare_benchmark.py
Original file line number Diff line number Diff line change
Expand Up @@ -36,7 +36,6 @@
import csv
import sys
from pathlib import Path
from typing import Dict, Tuple

METRIC = "us"
KERNEL_COLS = ("kernelName1", "kernelName2")
Expand Down Expand Up @@ -108,8 +107,8 @@ def _natural_key(val: str):
return (1, val)


Row = Dict[str, str]
Key = Tuple[Tuple[str, str], ...]
Row = dict[str, str]
Key = tuple[tuple[str, str], ...]


def _normalize_row(raw: Row) -> Row:
Expand All @@ -121,7 +120,7 @@ def _normalize_row(raw: Row) -> Row:
}


def _read_csv_rows(path: Path) -> Tuple[list[Row], Tuple[str, ...]]:
def _read_csv_rows(path: Path) -> tuple[list[Row], tuple[str, ...]]:
"""Return (rows, fieldnames). Whitespace stripped from values."""
if not path.exists():
raise SystemExit(f"input csv not found: {path}")
Expand All @@ -136,7 +135,7 @@ def _read_csv_rows(path: Path) -> Tuple[list[Row], Tuple[str, ...]]:
return rows, tuple(reader.fieldnames)


def _key_cols(base_cols: Tuple[str, ...], cur_cols: Tuple[str, ...]) -> Tuple[str, ...]:
def _key_cols(base_cols: tuple[str, ...], cur_cols: tuple[str, ...]) -> tuple[str, ...]:
"""Stable key column order across baseline/current schema drift."""
cols = []
seen = set()
Expand All @@ -148,8 +147,8 @@ def _key_cols(base_cols: Tuple[str, ...], cur_cols: Tuple[str, ...]) -> Tuple[st
return tuple(cols)


def _index_rows(rows: list[Row], key_cols: Tuple[str, ...]) -> Dict[Key, Row]:
indexed: Dict[Key, Row] = {}
def _index_rows(rows: list[Row], key_cols: tuple[str, ...]) -> dict[Key, Row]:
indexed: dict[Key, Row] = {}
for row in rows:
key = tuple(sorted((c, row.get(c, "")) for c in key_cols))
indexed[key] = row
Expand All @@ -167,7 +166,7 @@ def _parse_us(raw: Row) -> float | None:


def _fmt_key_compact(
key: Key, key_cols_order: Tuple[str, ...], constants: Dict[str, str]
key: Key, key_cols_order: tuple[str, ...], constants: dict[str, str]
) -> str:
"""Format key showing only cols whose value is NOT in `constants`."""
d = dict(key)
Expand All @@ -178,7 +177,7 @@ def _fmt_key_compact(
return " ".join(parts) if parts else "(common)"


def _find_constants(keys: list[Key], key_cols_order: Tuple[str, ...]) -> Dict[str, str]:
def _find_constants(keys: list[Key], key_cols_order: tuple[str, ...]) -> dict[str, str]:
"""Return cols whose value is identical across all `keys`."""
if not keys:
return {}
Expand Down Expand Up @@ -304,7 +303,7 @@ def _entry_sort_key(e):
if c not in HIDE_DISPLAY_COLS:
display_cols.append(c)

def _cell_value(c: str, d: Dict[str, str]) -> str:
def _cell_value(c: str, d: dict[str, str]) -> str:
if c in DERIVED_TUPLE_COLS:
srcs = DERIVED_TUPLE_COLS[c]
return "(" + ", ".join(d.get(s, "") for s in srcs) + ")"
Expand Down
6 changes: 5 additions & 1 deletion .github/workflows/pre-checks.yaml
Original file line number Diff line number Diff line change
Expand Up @@ -45,7 +45,11 @@ jobs:
with:
python-version: "3.12"
- name: Install dependencies
run: pip3 install ruff
# Pinned on purpose. Unpinned, a ruff release that widens the default
# rule set turns unrelated PRs red with no change on our side (0.16.0
# did exactly that). Bump this deliberately, with the resulting
# findings burned down in the same or a follow-up PR.
run: pip3 install ruff==0.16.0
- name: Install reviewdog
uses: reviewdog/action-setup@e04ffabe3898a0af8d0fb1af00c188831c4b5893
- name: Run ruff with reviewdog
Expand Down
110 changes: 60 additions & 50 deletions aiter/__init__.py
Original file line number Diff line number Diff line change
@@ -1,10 +1,11 @@
# SPDX-License-Identifier: MIT
# Copyright (C) 2024-2026, Advanced Micro Devices, Inc. All rights reserved.

import torch
import logging
import os
import sys
import logging

import torch

logger = logging.getLogger("aiter")

Expand Down Expand Up @@ -73,7 +74,7 @@ def getLogger():
if AITER_TRITON_ONLY:
logger.info("Triton ops only: CK and HIP ops (and their JIT build) are skipped.")
elif AITER_AOT_IMPORT:
from .jit import core as core # noqa: E402
from .jit import core as core
else:
# NOTE: do NOT wrap this block in try/except.
# Catching ImportError here silently truncates the top-level aiter
Expand All @@ -84,62 +85,71 @@ def getLogger():
# opus is gfx950-only but the package self-guards (warn + stubs on
# non-gfx950) inside aiter/ops/opus/__init__.py, so its import line
# is safe to put at top-level without try/except.
from .jit import core as core # noqa: E402
from .utility import dtypes as dtypes # noqa: E402
from .ops.enum import * # noqa: F403,E402
from .ops.norm import * # noqa: F403,E402
from .ops.quant import * # noqa: F403,E402
from .ops.gemm_op_a8w8 import * # noqa: F403,E402
from .ops.gemm_op_a16w16 import * # noqa: F403,E402
from .ops.gemm_op_a4w4 import * # noqa: F403,E402
from .ops.gemm_op_a8w4 import * # noqa: F403,E402
from .ops.batched_gemm_op_a8w8 import * # noqa: F403,E402
from .ops.batched_gemm_op_bf16 import * # noqa: F403,E402
from .ops.deepgemm import * # noqa: F403,E402
from .ops.opus import * # noqa: F403,E402
from .ops.aiter_operator import * # noqa: F403,E402
from .ops.activation import * # noqa: F403,E402
from .ops.attention import * # noqa: F403,E402
from .ops.custom import * # noqa: F403,E402
from .ops.custom_all_reduce import * # noqa: F403,E402
from .ops.quick_all_reduce import * # noqa: F403,E402
from .ops.moe_op import * # noqa: F403,E402
from .ops.moe_sorting import * # noqa: F403,E402
from .ops.moe_sorting_opus import * # noqa: F403,E402
from .ops.moe_mxfp4_aux import * # noqa: F403,E402
from .ops.pa_sparse_prefill_opus import * # noqa: F403,E402
from .ops.pos_encoding import * # noqa: F403,E402
from .ops.cache import * # noqa: F403,E402
from .ops.rmsnorm import * # noqa: F403,E402
from .ops.communication import * # noqa: F403,E402
from .ops.rope import * # noqa: F403,E402
from .ops.topk import * # noqa: F403,E402
from .ops.topk_plain import topk_plain # noqa: F403,F401,E402
from .ops.mha import * # noqa: F403,E402
from .ops.gradlib import * # noqa: F403,E402
from .ops.trans_ragged_layout import * # noqa: F403,E402
from .ops.sample import * # noqa: F403,E402
from .ops.fused_qk_norm_mrope_cache_quant import * # noqa: F403,E402
from .ops.fused_qknorm_idxrqknorm import ( # noqa: F401,E402
# isort: off
# Order below is load-bearing, do not sort. `dtypes` must be bound on the
# aiter package before any submodule that does `from aiter import dtypes`
# (mla.py among others) gets pulled in, otherwise that resolves against a
# partially initialised aiter and raises ImportError.
from .jit import core as core
from .utility import dtypes as dtypes
from .ops.enum import *
from .ops.norm import *
from .ops.quant import *
from .ops.gemm_op_a8w8 import *
from .ops.gemm_op_a16w16 import *
from .ops.gemm_op_a4w4 import *
from .ops.gemm_op_a8w4 import *
from .ops.batched_gemm_op_a8w8 import *
from .ops.batched_gemm_op_bf16 import *
from .ops.deepgemm import *
from .ops.opus import *
from .ops.aiter_operator import *
from .ops.activation import *
from .ops.attention import *
from .ops.custom import *
from .ops.custom_all_reduce import *
from .ops.quick_all_reduce import *
from .ops.moe_op import *
from .ops.moe_sorting import *
from .ops.moe_sorting_opus import *
from .ops.moe_mxfp4_aux import *
from .ops.pa_sparse_prefill_opus import *
from .ops.pos_encoding import *
from .ops.cache import *
from .ops.rmsnorm import *
from .ops.communication import *
from .ops.rope import *
from .ops.topk import *
from .ops.topk_plain import topk_plain # noqa: F401
from .ops.mha import *
from .ops.gradlib import *
from .ops.trans_ragged_layout import *
from .ops.sample import *
from .ops.fused_qk_norm_mrope_cache_quant import *
from .ops.fused_qknorm_idxrqknorm import ( # noqa: F401
fused_qknorm_idxrqknorm,
)
from .ops.fused_qk_norm_rope_cache_quant import * # noqa: F403,E402
from .ops.fused_qk_rmsnorm_group_quant import * # noqa: F403,E402
from .ops.groupnorm import * # noqa: F403,E402
from .ops.mhc import * # noqa: F403,E402
from .ops.causal_conv1d_update import * # noqa: F403,E402
from .ops.fused_split_gdr_update import * # noqa: F403,E402
from . import mla # noqa: F403,F401,E402
from .ops.fused_qk_norm_rope_cache_quant import *
from .ops.fused_qk_rmsnorm_group_quant import *
from .ops.groupnorm import *
from .ops.mhc import *
from .ops.causal_conv1d_update import *
from .ops.fused_split_gdr_update import *
from . import mla # noqa: F401

# isort: on

# Import Triton-based communication primitives from ops.triton.comms (optional, only if Iris is available)
try:
from .ops.triton.comms import (
IRIS_COMM_AVAILABLE,
IrisCommContext, # noqa: F401
calculate_heap_size, # noqa: F401
reduce_scatter as iris_reduce_scatter, # noqa: F401 # avoid shadowing C++ reduce_scatter exported by custom_all_reduce.py
all_gather, # noqa: F401
calculate_heap_size, # noqa: F401
reduce_scatter_rmsnorm_quant_all_gather, # noqa: F401
IRIS_COMM_AVAILABLE, # noqa: F401
)
from .ops.triton.comms import (
reduce_scatter as iris_reduce_scatter, # noqa: F401 # avoid shadowing C++ reduce_scatter exported by custom_all_reduce.py
)
except (ImportError, AttributeError):
# Iris or triton not available, skip import
Expand Down
5 changes: 3 additions & 2 deletions aiter/aot/asm_mla_decode_fwd.py
Original file line number Diff line number Diff line change
@@ -1,6 +1,7 @@
from collections import namedtuple
import os
import concurrent.futures
import os
from collections import namedtuple

from csrc.cpp_itfs.mla.asm_mla_decode_fwd import compile
from csrc.cpp_itfs.utils import AITER_CORE_DIR

Expand Down
11 changes: 6 additions & 5 deletions aiter/aot/flydsl/common.py
Original file line number Diff line number Diff line change
Expand Up @@ -5,17 +5,18 @@

from __future__ import annotations

from contextlib import contextmanager
import enum
import json
import multiprocessing
from multiprocessing.connection import wait as wait_for_sentinels
import os
import shutil
import tempfile
import time
from collections.abc import Callable, Iterator
from contextlib import contextmanager
from dataclasses import dataclass
import enum
import os
from typing import Any, Callable, Iterator
from multiprocessing.connection import wait as wait_for_sentinels
from typing import Any

_DEFAULT_KERNEL_TIMEOUT = 1200.0
_DEFAULT_MAX_WORKERS = 64
Expand Down
5 changes: 2 additions & 3 deletions aiter/aot/flydsl/gemm.py
Original file line number Diff line number Diff line change
Expand Up @@ -34,7 +34,6 @@
import re
import sys
import time
from typing import Dict, Optional

import flydsl.expr as fx

Expand Down Expand Up @@ -83,7 +82,7 @@
}


def _parse_bool(value: Optional[str]) -> bool:
def _parse_bool(value: str | None) -> bool:
if value is None:
return False
normalized = value.strip().lower()
Expand All @@ -96,7 +95,7 @@ def _parse_bool(value: Optional[str]) -> bool:
raise ValueError(f"Expected True/False, got {value!r}")


def _parse_preshuffle_kernel_name(name: str) -> Optional[Dict]:
def _parse_preshuffle_kernel_name(name: str) -> dict | None:
m = _PRESHUFFLE_RE.fullmatch(name)
if m is None:
return None
Expand Down
34 changes: 17 additions & 17 deletions aiter/aot/flydsl/grouped_moe.py
Original file line number Diff line number Diff line change
Expand Up @@ -439,23 +439,23 @@ def compile_one_config(**job):
)
warp_tile_m = job["tile_m"] // job["m_warp"]
contiguous = bool(job.get("grouped_contiguous_m", False))
common = dict(
model_dim=job["model_dim"],
inter_dim=job["inter_dim"],
experts=job["experts"],
max_m=job["max_m"],
tile_m=job["tile_m"],
tile_n=job["tile_n"],
tile_k=job["tile_k"],
m_warp=job["m_warp"],
n_warp=job["n_warp"],
out_dtype=job["out_dtype"],
num_buffers=job["num_buffers"],
grouped_persistent_m=job["grouped_persistent_m"],
grouped_contiguous_m=contiguous,
persistent_workers=job["persistent_workers"],
expert_sched_mode=job["expert_sched_mode"],
)
common = {
"model_dim": job["model_dim"],
"inter_dim": job["inter_dim"],
"experts": job["experts"],
"max_m": job["max_m"],
"tile_m": job["tile_m"],
"tile_n": job["tile_n"],
"tile_k": job["tile_k"],
"m_warp": job["m_warp"],
"n_warp": job["n_warp"],
"out_dtype": job["out_dtype"],
"num_buffers": job["num_buffers"],
"grouped_persistent_m": job["grouped_persistent_m"],
"grouped_contiguous_m": contiguous,
"persistent_workers": job["persistent_workers"],
"expert_sched_mode": job["expert_sched_mode"],
}
if contiguous:
act_lead = 1
ub = job["token_num"] * job["topk"] + job["experts"] * (job["tile_m"] - 1)
Expand Down
3 changes: 1 addition & 2 deletions aiter/aot/flydsl/moe.py
Original file line number Diff line number Diff line change
Expand Up @@ -27,7 +27,6 @@
import os
import sys
import time
from typing import Optional

from aiter.aot.flydsl.common import (
collect_aot_jobs,
Expand Down Expand Up @@ -202,7 +201,7 @@ def _precompile_to_cache(
# pin ``waves_per_eu`` in ``get_flydsl_stage{1,2}_kernels`` (only the
# production-variant ``_persist_async_w4_cumul3`` does), causing
# ``AOT cache miss`` at runtime even though the .pkl is present on disk.
waves_per_eu: Optional[int] = None,
waves_per_eu: int | None = None,
k_batch: int = 1,
b_nt: int = 2,
gate_mode: str = "separated",
Expand Down
2 changes: 1 addition & 1 deletion aiter/aot/flydsl/mxfp4_moe.py
Original file line number Diff line number Diff line change
Expand Up @@ -57,12 +57,12 @@ def _job_key(job: dict) -> tuple:

def parse_csv(csv_path: str):
"""Parse an fp4 tuned CSV into unique mxmoe-port compile jobs (one per stage)."""
from aiter.ops.flydsl.mxfp4_gemm2_kernels import _epilog_of
from aiter.ops.flydsl.mxfp4_kname import (
_is_mxfp4_kname,
_parse_mxfp4_g1_kname,
_parse_mxfp4_g2_kname,
)
from aiter.ops.flydsl.mxfp4_gemm2_kernels import _epilog_of

jobs = []
seen = set()
Expand Down
5 changes: 3 additions & 2 deletions aiter/aot/pa.py
Original file line number Diff line number Diff line change
@@ -1,6 +1,7 @@
from collections import namedtuple
import os
import concurrent.futures
import os
from collections import namedtuple

from csrc.cpp_itfs.pa.pa import compile

PAConfig = namedtuple(
Expand Down
5 changes: 3 additions & 2 deletions aiter/aot/pa_ragged.py
Original file line number Diff line number Diff line change
@@ -1,6 +1,7 @@
from collections import namedtuple
import os
import concurrent.futures
import os
from collections import namedtuple

from csrc.cpp_itfs.pa.pa_ragged import compile

PAConfig = namedtuple(
Expand Down
5 changes: 3 additions & 2 deletions aiter/aot/pa_v1.py
Original file line number Diff line number Diff line change
@@ -1,6 +1,7 @@
from collections import namedtuple
import os
import concurrent.futures
import os
from collections import namedtuple

from csrc.cpp_itfs.pa.pa_v1 import compile

PAConfig = namedtuple(
Expand Down
Loading
Loading