[Triton/Gluon] [ASM] [HIP] Add Kimi-K3 KDA decode kernels on gfx950

.github/instructions/aiter-ops-triton.instructions.md

@@ -45,9 +45,10 @@ that duplicates functionality already in the tree, even partially — the fix is
 to extend or import the existing implementation, not to add a parallel copy.
 
 `utils/` is layered on purpose: `config_utils.py` holds the shared core
-(`resolve_config_dir`, `load_config_json`, the path constants) and each family
-keeps its own loader module (`gemm_config_utils`, `conv_config_utils`,
-`mhc_config_utils`, `moe_config_utils`, `tuned_config_utils`) on top of it.
+(`resolve_config_dir`, `load_config_json`, `select_leq_config`, the path
+constants) and each family keeps its own loader module (`gemm_config_utils`,
+`conv_config_utils`, `mhc_config_utils`, `moe_config_utils`,
+`quant_config_utils`, `tuned_config_utils`) on top of it.
 Flag a function given a second home — a re-export, a wrapper that only
 forwards to another module, or a copy of a core helper inside a family module.
 
@@ -69,6 +70,19 @@ and tuned JSON in `configs/`. Flag:
   Gluon bodies under `_gluon_kernels/<arch>/`, mirroring the wrapper's
   category path. JIT-decorated device helpers that are called only from
   another kernel may remain with the entry kernel they support.
+- A kernel module or directory whose path under `_triton_kernels/`
+  (`_gluon_kernels/<arch>/`) differs from the folder of the wrapper that
+  launches it. A kernel directory is named after the wrapper folder, never
+  after the kernel family or a topic: `normalization/`, not `norm/`. One
+  existing directory predates this rule and is grandfathered until a
+  dedicated follow-up moves it and all its importers: the vendored
+  `_triton_kernels/flash_attn_triton_amd/` (launched by `attention/mha.py`).
+  Do not flag code in it; flag any new mismatched directory. A kernel that
+  wrappers in several folders reuse lives with the wrapper that owns it, or in
+  `_triton_kernels/common/` when none does (`common/splitk_reduce.py`);
+  importing it across folders is fine. A PR that moves a wrapper into another
+  folder moves its kernels too and updates every importer of the old module
+  path (`grep -rn "_triton_kernels.<old>"`).
 - Generic helpers (config loading, shuffling, arch detection, logging)
   re-implemented inside a kernel file instead of imported from `utils/`.
 - New code importing via the legacy flat paths
@@ -104,7 +118,7 @@ their tuned configs can be imported by a framework that is not PyTorch
 Every tuned config lives in one nested layout:
 `configs/<arch>/<backend>/<op>/<d_type>/`, e.g.
 `configs/gfx950/triton/gemm/gemm_afp4wfp4/DEFAULT.json`. `<op>` is `gemm`,
-`moe`, `conv`, `mhc`, `attention`, `gmm` or `fusions`; `<d_type>` is
+`moe`, `conv`, `mhc`, `attention`, `gmm`, `fusions` or `quant`; `<d_type>` is
 `config_name.lower().replace("-", "_")`. The flat arch-prefixed directories
 and every fallback that reached them are gone. Flag:
 
@@ -215,14 +229,19 @@ values for either backend live in JSON, never in Python. Flag:
   family loader or `resolve_config_dir()` would work — a hand-built path is a
   second place the layout is encoded, and it skips the argument validation
   that makes a wrong value fail closed.
+- A hand-written loop selecting the smallest matching `N_LEQ_*` (or another
+  upper-bound prefix) entry, or a hand-written multi-axis bucket walk — use
+  `select_leq_config()` (`axes=` for several axes) so threshold ordering,
+  fallback, and copying semantics have one implementation.
 - A second MOE config reader. `utils/moe_config_utils.py::get_moe_dispatch` is
   the only MOE fetcher; flag any new MOE path built by hand, any direct
   `load_config_json` on a `moe/` file, and any reintroduced per-wrapper MOE
   loader.
 - A new arch- or backend-fallback chain inside a loader (try this arch, then
-  that one; try triton, then gluon). Resolution is deterministic. MHC's gfx942
-  fallback is the one documented exception and it goes through the `arch=`
-  override, not through a probe.
+  that one; try triton, then gluon). Resolution is deterministic. The
+  documented compatibility exceptions are MHC's gfx942 fallback and Triton
+  `fused_clamp_act_mul`'s legacy gfx950 fallback; both use the `arch=` override
+  instead of a probe.
 - A raw config list handed to `@triton.autotune`. Route it through
   `autotune_configs` from `aiter.ops.triton.utils.tuned_config_utils`:
 
@@ -349,6 +368,31 @@ All weight/scale pre-shuffle helpers are unified in
   inside the category folder, shared helpers in the existing
   `*_test_utils.py` / `utils/` modules. Flag tests added flat at the
   `op_tests/triton_tests/` root or as one-off scripts.
+- A test's folder is the folder of the wrapper it imports. The path under
+  `op_tests/triton_tests/` mirrors the wrapper's path under
+  `aiter/ops/triton/`: a test of `aiter/ops/triton/attention/mla.py` is
+  `op_tests/triton_tests/attention/test_mla.py`, a test of
+  `gemm/basic/gemm_a8w8.py` is `gemm/basic/test_gemm_a8w8.py`, a test of
+  `gated_delta_net/fused_kda_decode.py` is
+  `gated_delta_net/test_fused_kda_decode.py`. The wrapper's folder decides
+  the test folder, not the kernel directory: `attention/mha.py` keeps its
+  tests in `triton_tests/attention/` even though its kernels sit in the
+  vendored `_triton_kernels/flash_attn_triton_amd/`; a sub-package counts as
+  a folder (`moe/moe_routing/routing.py` → `triton_tests/moe/moe_routing/`).
+  The few wrappers that still sit flat at `aiter/ops/triton/<op>.py`
+  (`activation.py`, `topk.py`, ...) keep their tests flat at the
+  `triton_tests/` root; `utils/` helpers are tested under
+  `triton_tests/utils/`; `torch_compile/` and `triton_metadata_redirect/`
+  are infrastructure suites, not op folders. Flag:
+  - A test whose folder differs from the folder of the wrapper it imports
+    (`from aiter.ops.triton.gated_delta_net...` in a test under
+    `triton_tests/attention/` or at the root).
+  - A test placed by kernel directory or by topic instead of by wrapper
+    folder, and a new test folder that does not match a folder under
+    `aiter/ops/triton/` (every op test folder also carries an `__init__.py`).
+  - A PR that moves a wrapper into another folder without moving its kernels
+    and its test, or that moves a test without updating every importer of the old module
+    path (`grep -rn "triton_tests.<old>"`).
 - No kernel tuning configs in test files: flag test code that hardcodes
   config dicts (`BLOCK_SIZE_*`, `num_warps`, `waves_per_eu`, ...) or passes
   literal `config=` overrides to a wrapper. Tests exercise the wrapper's own

.github/scripts/prebuild_atom_kimi_k3.py

@@ -0,0 +1,53 @@
+#!/usr/bin/env python3
+# SPDX-License-Identifier: MIT
+# Copyright (C) 2026, Advanced Micro Devices, Inc. All rights reserved.
+
+"""Build Kimi-K3's slow GEMM dependency before the ATOM startup deadline."""
+
+import importlib
+import sys
+from pathlib import Path
+
+MODULE = "module_gemm_a8w8_bpreshuffle_cktile"
+EXPECTED_ARCH = "gfx950"
+
+
+def main() -> None:
+    root = Path(__file__).resolve().parents[2]
+    # Use the build-time import path without initializing AITER's runtime.
+    sys.path.insert(0, str(root / "aiter"))
+    core = importlib.import_module("jit.core")
+
+    args = core.get_args_of_build(MODULE)
+    core.build_module(
+        md_name=MODULE,
+        srcs=args["srcs"],
+        flags_extra_cc=args["flags_extra_cc"],
+        flags_extra_hip=args["flags_extra_hip"],
+        blob_gen_cmd=args["blob_gen_cmd"],
+        extra_include=args["extra_include"],
+        extra_ldflags=args["extra_ldflags"],
+        verbose=args["verbose"],
+        is_python_module=args["is_python_module"],
+        is_standalone=args["is_standalone"],
+        torch_exclude=args["torch_exclude"],
+        third_party=args.get("third_party", []),
+        hipify=args.get("hipify", False),
+        flags_extra_hip_per_source=args.get("flags_extra_hip_per_source", {}),
+    )
+
+    target = Path(core.get_user_jit_dir()) / f"{MODULE}.so"
+    expected = root / "aiter" / "jit" / target.name
+    if target.resolve() != expected.resolve() or not target.is_file():
+        raise RuntimeError(f"Prebuild did not create {expected}: got {target}")
+
+    arches = core._so_offload_archs(target)
+    if arches != {EXPECTED_ARCH}:
+        raise RuntimeError(f"Unexpected GPU architectures: {sorted(arches)}")
+    print(
+        f"Prebuilt {MODULE}: {target.stat().st_size} bytes, {EXPECTED_ARCH}", flush=True
+    )
+
+
+if __name__ == "__main__":
+    main()

.github/scripts/split_tests.sh

@@ -334,6 +334,7 @@ elif [[ "$TEST_TYPE" == "triton" ]]; then
     FILE_TIMES[op_tests/triton_tests/moe/test_moe_routing_herd.py]=14
     FILE_TIMES[op_tests/triton_tests/test_fused_rearrange_sigmoid_gdr.py]=14
     FILE_TIMES[op_tests/triton_tests/attention/test_pa_prefill_sparse.py]=12
+    FILE_TIMES[op_tests/triton_tests/gated_delta_net/test_fused_kda_decode.py]=12
     FILE_TIMES[op_tests/triton_tests/gemm/basic/test_gemm_a16wfp4.py]=12
     FILE_TIMES[op_tests/triton_tests/gemm/basic/test_gemm_a8wfp4.py]=12
     FILE_TIMES[op_tests/triton_tests/chunk_delta_attn/test_fast_launch.py]=10
@@ -357,7 +358,6 @@ elif [[ "$TEST_TYPE" == "triton" ]]; then
     FILE_TIMES[op_tests/triton_tests/torch_compile/test_compile_activation.py]=5
     FILE_TIMES[op_tests/triton_tests/attention/test_sparse_attention_dsv4_bwd.py]=4
     FILE_TIMES[op_tests/triton_tests/fusions/test_fused_rmsnorm_indexed_adaln.py]=4
-    FILE_TIMES[op_tests/triton_tests/test_fused_kda_decode.py]=4
     FILE_TIMES[op_tests/triton_tests/test_softmax.py]=4
     FILE_TIMES[op_tests/triton_tests/torch_compile/test_compile_fused_mul_add.py]=4
     FILE_TIMES[op_tests/triton_tests/torch_compile/test_compile_quant_per_token.py]=4

.github/workflows/actionlint.yaml

@@ -26,7 +26,7 @@ jobs:
 
       - name: Install actionlint
         env:
-          ACTIONLINT_VERSION: "1.7.7"
+          ACTIONLINT_VERSION: "1.7.12"
         run: |
           set -euo pipefail
           curl -fsSL \