Skip to content

[HIP] [CK] [JIT] [Tuner] Rank accuracy-qualified kernels by package energy - #6970

Merged
yzhou103 merged 14 commits into
ROCm:mainfrom
rpoornac:power-kernel-select
Nov 3, 2026
Merged

yzhou103 merged 14 commits into
ROCm:mainfrom
rpoornac:power-kernel-select

Conversation

@rpoornac

@rpoornac rpoornac commented Oct 26, 2026 •

Copy link
Copy Markdown
Contributor

Summary

  • Add --select power (experimental, opt-in) so tuners that support it keep the kernel with the lowest package energy (us × PPT watts) among kernels that pass errRatio. Latency remains the default, and a latency tune writes the same CSV schema as before.
  • Every latency-or-power decision goes through one SelectionPolicy (aiter/utility/selection_policy.py), built from --select. Every power read goes through one function, package_power.measure_kernel. Tuners do not branch on the select mode and do not read or write AITER_SELECT. AITER_SELECT=power only makes runtime dispatch overlay tables from aiter/configs/power_configs/.
  • Power is the mean power1_input (PPT) of the HIP device under test while the timed kernel launches back to back for AITER_POWER_WINDOW_S (default 0.4 s). When host launch gaps leave the GPU partly idle, the reading is scaled to active power from the launch count and the timed us. Windows under 20% busy stay unmeasured.
  • Power rows record select=power, perf_level, power_cap_w, and sclk_max_mhz, and a power tune of a shipped table is written to aiter/configs/power_configs/. Only rows whose watts were measured are written; unmeasured shapes stay untuned for the next power tune, and the AITER_SELECT=power overlay uses only rows stamped select=power.
  • Latency tunes keep main's output schemas, including the MoE -o2 profile and profile_fmoe.csv.

PR Self-Review

Reviewed revision

bbc0ae5

Behavior and intent

Tuners rank candidates by latency only, so the tuned tables always pick the fastest kernel even when a slightly slower one uses much less package energy. This PR adds an opt-in --select power mode that ranks accuracy-qualified candidates by us × package watts and stores those rows separately (aiter/configs/power_configs/), loaded only with AITER_SELECT=power. Default behavior is unchanged: --select latency keeps the lowest us and writes no power columns. Supported tuners: the shared GEMM tuners (GemmCommonTuner: a16w16 incl. hipBLASLt/asm/FlyDSL/Opus/Triton/torch, a8w8, blockscale, gradlib), the fused-MoE CK two-stage, grouped, and FlyDSL MXFP4 paths, the GDN K5 BV sweep, and IQ2R. Conv3d, a6w6, and the mixed-MXFP tuners (a6w4/a4w6) reject --select power, because their gain guard keeps a kernel only for speed.

Validation and unverified scope

Hardware: 8× MI355X (gfx950), ROCm 7.2.4, power cap 1400 W, clocks auto (SCLK 500-2400 MHz), GPUs otherwise idle.

  • Unit: python -m unittest op_tests.tuning_tests.test_fast_tuning — 54 tests pass, including a real-sensor test (8192³ bf16 GEMM must read > 1.5× idle and ≥ 0.85× a back-to-back loop; skips without power1_input).
  • Other tuning suites pass: test_mp_tuner_logic, test_compare_logic, test_config_shape_collision, test_fhmoe_tuner, test_a6w6_tuning, test_mixed_mxfp_tuning. test_tuner_infra has 2 failures (a_preshuffle blank-cell) that also occur on main ([FlyDSL] Add gfx1250 grouped MoE A-preshuffle and tuned CSV validation #6709).
  • End to end, latency and power tunes on GPUs 0-7, low/mid/high shapes from the shipped untuned CSVs: a16w16 (all backends + hipBLASLt, 4 shapes), gradlib (4), a8w8 (3), a8w8 blockscale (3), CK MoE ASM FP8 (3), MXFP4 staged (3) and fast scan (2), IQ2R (3), and a16w16 with --mp 2. On every shape the latency CSV has no power columns, every power row has finite watts and the conditions, and the power pick is the minimum-energy measured candidate. Examples: a8w8 16384×8192×1024 is 23% slower and 9% lower energy; blockscale 1024×1536×7168 is 24% slower and 16% lower energy; MXFP4 32 tokens is 4% slower and 18% lower energy.
  • MoE profile schema (FP8 per_1x128, ASM, 3 shapes): the latency -o2 profile has main's 29 columns with no power columns, and profile_fmoe.csv has none; the power run keeps the energy-ranked profile and stamps all rows.
  • Power sample vs CUDA-graph replay (no launch gaps): torch.mm within ~10%; split-K FlyDSL still reads 12-28% low after the launch-gap correction, though the energy ranking on that shape matches the graph-replay ranking.
  • Unverified: grouped MoE (gfx1250 only) is covered by unit tests only; K5 has no shipped untuned shapes; gfx942 not run; kernels under ~5 µs are left unmeasured by design; the power score is an isolated replay, so a winner can change in a model, at another power cap, or next to other kernels; end-to-end model perf/W (MLPerf) is not part of this PR. Graph-replay power sampling is planned as a follow-up.

PR-specific question

No kernel dispatch path changes. Which kernel runs is still chosen by the tuned CSV for the arch/dtype/shape keys, as before. A power row can reach dispatch only through AITER_SELECT=power (overlay of aiter/configs/power_configs/) or an explicit AITER_CONFIG_* path. Without AITER_SELECT=power, power_configs/ is ignored. A power row is written only for a candidate that passed the same errRatio accuracy gate as latency mode and has finite measured watts; a batch with no finite watts stays a latency result. Merging a power row and a latency row for one shape, or power rows tuned under different perf level, cap, or SCLK, raises instead of auto-resolving. Tuners that cannot rank by energy reject --select power (SUPPORTS_POWER_SELECT = False).

Changed op tests

op_tests/tuning_tests/test_fast_tuning.py — the only changed op test. Its new cases cover power selection, and none duplicate an existing case:

  • Ranking and policy: energy vs latency order, unmeasured rows ranked last, MoE stage-energy pairs, candidate_beats, SelectionPolicy decisions, mp_tuner using the activated policy, a16w16 passing its policy to hipBLASLt.
  • Measurement: the timed kernel is replayed without run_perftest options; bounded in-flight launches; launch-gap correction; real-sensor large-GEMM check.
  • Output and loading: only measured rows are kept and stamped, power output redirected to power_configs/, latency frames and the latency MoE profile without power columns, power/latency merge conflicts, the power_configs/ overlay (unstamped rows and files without select ignored).
  • Build safety: jit.core loaded as in setup.py (with and without duplicate shapes) does not import the aiter package.
  • Guard: no tuner file decides the select mode, reads power, or uses AITER_SELECT itself.

Test plan

  • python -m unittest op_tests.tuning_tests.test_fast_tuning (54 tests, MI355X)
  • Latency and power tunes on 8 tuner paths, low/mid/high shapes, MI355X (see Validation above)
  • MLPerf and other end-to-end perf/W runs (out of scope for this PR)

Made with Cursor

@rpoornac
rpoornac requested a review from a team October 26, 2026 23:11
@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:gfx1250-ffm-triton Run the five-shard gfx1250 FFM Triton test suite
ci:triton-300x Run an additional Triton test job on MI300X in PRs (added automatically when gfx942 configs change); main branch always runs both MI35X and MI300X
ci:triton-355 Run the full Triton test suite on MI35X, not only the tests the change affects
ci:windows Windows JIT op smoke test on the gfx1201 ROCm runner
ci:multi-gpu Aiter multi-GPU tests on the 8-GPU runner
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: DeepSeek-V4-Pro, MiniMax-M3-MXFP4
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 6970 --add-label <label>

One backend per PR:
A PR changes one kernel backend: [Triton/Gluon] (Triton and Gluon count as one), [HIP], [ASM], [CK], [OPUS] or [FlyDSL]. If the title ends up with two backend tags, split the PR -- as stacked pull requests when one part cannot merge without the other.

PR title tags & labels:
Component tags ([Triton/Gluon], [HIP], [CK], [ASM], ...) are added to the PR title and as PR labels automatically from the changed files and re-synced on every push — change-type tags like [fix]/[Perf], op tags like [MLA], and human labels (ci:*) are left untouched. Add the no-auto-title label to stop the title rewrites; labels stay in sync either way.

@github-actions github-actions Bot changed the title [Tuner] Rank accuracy-qualified kernels by package energy [HIP] [CK] [JIT] [Tuner] Rank accuracy-qualified kernels by package energy Oct 26, 2026
@zufayu
zufayu requested a review from yzhou103 October 27, 2026 01:11
@yzhou103
yzhou103 requested a review from zufayu October 27, 2026 03:23
Comment thread aiter/utility/mp_tuner.py Resolved
@yzhou103

Copy link
Copy Markdown
Contributor

Thanks for putting this together. Power-aware kernel selection is an interesting and potentially valuable direction, especially for sustained workloads where energy efficiency is an important objective.
But I have a design concern about whether the current measurement can provide a stable, workload-independent power score.
The direction is valuable, but the current implementation may not produce a stable, workload-independent power score.

The tuner measures package PPT while replaying each kernel in isolation for approximately 0.4 seconds, under whatever clock, thermal, and power-management policy the firmware currently applies. This may not represent the kernel’s behavior inside a real model:

Small shapes may be dominated by telemetry granularity, launch overhead, or background power.
The replay itself may push the GPU into a sustained power cap.
Candidates are measured sequentially without cooldown or randomized ordering, so one candidate may affect the thermal/frequency state of the next one.
The same shape may behave differently across models. A kernel selected by this tuner could improve perf/W in model A but be slower or less efficient in model B, because neighboring kernels, concurrency, communication, and the GPU’s thermal/power state change the operating point. Since the tuned CSV is primarily keyed by shape and reused across workloads, an isolated replay may select a winner that is not optimal for every model.
There is also no explicit latency or throughput constraint in the current objective:

score = runtime_us × package_watts
A substantially slower kernel can win as long as its measured power decreases proportionally. That may be valid for raw energy minimization, but it may violate production latency or throughput requirements.

Another concern is ownership of the problem. If tuning is performed at maximum frequency/power but the selected kernel performs poorly in production because the GPU reaches a power cap, are we observing a firmware or power-management issue? In that case, selecting a slower, lower-power kernel may hide the underlying problem instead of fixing it. The power cap and frequency policy may need to be addressed by firmware or the platform power manager, and the tuner should expose the issue rather than encode a workaround into the kernel configuration.

Before treating this as a production tuning strategy, could we clarify:

Is the goal to minimize energy subject to a latency/throughput bound, or to minimize raw runtime × watts?
How do we distinguish a genuinely more energy-efficient kernel from one that is simply being throttled differently by firmware?
Should this be validated under fixed clocks and explicit power-cap settings?
Do we have representative end-to-end results showing that the power-selected kernel improves perf/W without violating latency targets?
How do we handle cases where the best kernel differs between model A and model B?
I suggest keeping this feature explicitly experimental unless it adds a latency guard or Pareto-based selection, records the relevant clock/power-cap state, and is validated on representative model workloads. Otherwise, the tuner may mask a firmware power-management problem by selecting a power-friendly but slower operator.

@rpoornac

Copy link
Copy Markdown
Contributor Author

Thanks for the design notes. This PR keeps power selection opt-in. Latency is still the default, and the README now says so.

Answers to the questions, as the code stands:

  • The objective is raw package energy, runtime × PPT watts, among kernels that already pass errRatio. There is no latency or throughput bound on top of that. A slower kernel can win when its joules are lower.
  • The sample is an isolated replay of about 0.4 s on power1_input. It is not a model-level score. Sequential candidates can see different clocks, and the same shape can prefer a different kernel in another model because the CSV key is the shape.
  • We have not shown an end-to-end perf/watt win. At an open cap the package watts of different tiles sat together; the watt gap showed up with the clock held fixed. That is why this should not be the production default yet.
  • A missing sensor no longer changes the latency order, and a single AITER_POWER_PATH is rejected for a multi-process tune (see the two threads above).

A latency guard or a Pareto pick would be a follow-up. I would rather not add one in this PR until there is a fixed-clock or capped end-to-end run to set the bound against.

@yzhou103

Copy link
Copy Markdown
Contributor

Thanks, that clarifies that the feature is opt-in, is intended to use per-model tuned CSVs, and does not directly modify PM firmware behavior.

Before treating this as a product-facing feature, I think we still need to clarify the actual motivation and target scenario. What concrete workload or deployment problem is this intended to solve? For example, is there a real model running under a power cap where this selection provides a measurable benefit?

So far, the evidence appears to be based on isolated kernel replay. We do not yet have end-to-end model results showing improved perf/W, lower energy, or acceptable latency impact. Without that validation, it is difficult to determine whether this is ready for product release or should remain an experimental tuning option.

Also, AITER currently merges tuned CSVs by shape-based keys by default. How is model separation guaranteed in the actual config loading and merge path? If configs from different models are merged, model-specific power winners can still collapse into a single row and be applied across workloads.

Although this feature does not directly change PM firmware, it still selects kernels based on behavior observed under the current PM policy. That behavior may depend on power caps, clocks, thermal state, and the surrounding model workload.

@rpoornac

Copy link
Copy Markdown
Contributor Author

Thanks. Following up on my earlier note: we have end-to-end workload runs showing perf/W gains at a fixed power cap and at a fixed clock cap. I'd prefer to walk through that data with you offline rather than post it here. Until it is validated more widely, the feature stays experimental and opt-in.

On model separation, you're right that the default load merges tables by shape key. d9cd30c keeps power rows in aiter/configs/power_configs/. They load only with AITER_SELECT=power and replace only the shapes they cover. An explicit AITER_CONFIG_* path loads one model's table with no merge. A merge that finds a power row and a latency row for one shape, or power rows tuned under different conditions, now raises instead of auto-picking.

On power-management dependence: each power row now records the perf level, power cap, and SCLK range it was tuned under, so a winner from one state is not silently compared with one from another.

@yzhou103

Copy link
Copy Markdown
Contributor

I ran the latest head (94d257f) on an MI355X. The unit tests pass, and the raw PPT read works (about 310–326 W for a tiny GEMM, about 1372–1398 W for an 8k bf16 GEMM at the 1400 W cap). But an a16w16 tune with --select power --shape_grouped --mp 1 (torch and triton backends) writes watts and tflops_per_w as Null for every shape, while the rows are still stamped select=power.

Cause: _replay in mp_tuner.worker calls func(*args, **kwargs), but those kwargs are the run_perftest kwargs (num_warmup, num_iters, ...), so every replay raises TypeError: run_torch_gemm_a16w16() got an unexpected keyword argument 'num_warmup'. record_package_watts swallows the error and returns NaN, and the missing-watts fallback then keeps latency order. As a result, power mode silently produces latency picks on what looks like every mp_tuner-based tuner that passes perftest kwargs.

Comment thread csrc/kernels/iq2r/iq2r_glm53_tune.py Resolved
Comment thread aiter/utility/base_tuner.py Resolved
Comment thread aiter/utility/base_tuner.py Resolved
Comment thread aiter/utility/base_tuner.py Resolved
@rpoornac
rpoornac force-pushed the power-kernel-select branch from bfb22bb to 454b664 Compare October 29, 2026 02:42
@yzhou103

Copy link
Copy Markdown
Contributor

Power measurement on the mp_tuner path

Thanks for fixing the num_warmup TypeError; watts are now recorded. However, the new replay mostly measures idle time, so the watts can't be used for ranking yet.

Problem. replay_timed_call runs the full run_perftest(...) repeatedly inside the 0.4 s power window. Each call includes profiler start/stop, warmup, synchronization and trace parsing, so the GPU is idle most of the time. Same kernel on an MI355X (gfx950):

shape (bf16 mm)    us      back-to-back W   GPU busy   run_perftest replay W   GPU busy
64x7168x2048       13.8    541-638          ~100%      343-388                 ~3%
256x4096x4096      22.5    1067-1242        ~100%      367-553                 ~7%
4096^3             107.6   1340-1397        ~100%      576-873                 ~21%

A real gemm_a16w16_tune.py --select power --mp 1 run wrote 348 / 515 / 335 W for these shapes. These values are idle-dominated, like the replay column, and far below the back-to-back numbers. Shorter kernels are diluted more, so us * watts ends up close to latency order plus noise. The unit tests mock the measurement and don't catch this.

Scope. This affects mp_tuner.worker, i.e. every GemmCommonTuner subclass and gradlib. MoE, gdn_k5 and iq2r measure closures that launch the kernel directly, so they look fine. I checked those by reading the code, not by measuring.

Clarification. By "measure in the timing call" I meant sampling inside run_perftest's replay loop, not calling run_perftest again inside the window. Sorry for the ambiguity. Either fix is fine:

  1. Sample in the existing perftest replay hook (the AITER_SMI_MONITOR path). It replays func back to back with the same rotated args or graph as the timer. Note that it reads amdsmi current_socket_power rather than hwmon power1_input; please confirm the two agree, or use one sampler for both.

  2. Replay the kernel directly, with the run_perftest options stripped. This is the original _replay without the kwargs bug. I checked on gfx950 that it works and keeps the GPU fully busy:

    _PERFTEST_KEYS = set(inspect.signature(run_perftest).parameters) - {"func", "args", "kwargs"}
    
    def kernel_replay(func, args, kwargs):
        kernel_kwargs = {k: v for k, v in kwargs.items() if k not in _PERFTEST_KEYS}
        return lambda: func(*args, **kernel_kwargs)
    
    watts = record_package_watts(kernel_replay(func, args, kwargs), device_index=gpu_id)

Please also add one non-mocked test, e.g. watts for a large GEMM should be well above idle and close to a back-to-back loop. It would have caught both the NaN bug and this dilution.

@yzhou103

Copy link
Copy Markdown
Contributor
  • Minor: the prebuild guard doesn't cover the duplicate-shape branch. test_prebuild_config_load_does_not_import_aiter_package only exercises tables without duplicates; _raise_on_duplicate_shapes returns early there. With an overlapping shape (e.g. two AITER_CONFIG_BF16_BATCHED_GEMM files sharing one row), loading jit.core the way setup.py does fails with KeyError: 'module_aiter_core' in both latency and power mode, because the auto-dedup branch now imports sort_frame_for_selection / power_selection_enabled from aiter.utility. Main auto-resolves the same case. The shipped tables don't hit this, so it's an edge case, but a duplicate-shape case in that test would make the guard complete.

@rpoornac

Copy link
Copy Markdown
Contributor Author

Thanks @yzhou103. 5355cd8 covers both comments from this morning and the three maintainability points you sent separately.

Power sample measured idle time. I went with your option 2. package_power.kernel_call(func, args, kwargs) removes the run_perftest options, and measure_kernel launches the kernel back to back for the window. That is now the only path to the sensor. MI355X (gfx950), bf16 torch.mm, GPU 0 otherwise idle:

shape us old run_perftest replay W new sample W back-to-back loop W
64x7168x2048 16.1 356 697 806
256x4096x4096 21.8 478 983 1042
4096^3 116.2 601 1247 1279

The old column reproduces your table. New test test_power_sample_of_a_large_gemm_tracks_a_busy_gpu reads the real sensor and does not mock it. It runs an 8192^3 bf16 GEMM through the tuner's sample and requires more than 1.5x idle and at least 0.85x a back-to-back loop. Locally it read idle 238 W, loop 1313 W, tuner 1339 W. It skips when power1_input is not readable.

Prebuild guard missed the duplicate-shape branch. The dedup in _raise_on_duplicate_shapes now imports from aiter.utility only when a running import aiter asked for power tables. Otherwise it sorts by lowest us, as main does. test_prebuild_duplicate_shapes_resolve_without_aiter_package loads jit.core the way setup.py does, with two AITER_CONFIG_BF16_BATCHED_GEMM files sharing a row, in both latency and power mode. It fails on the previous commit with KeyError: 'module_aiter_core'.

Maintainability.

  1. Ranking in one place. New aiter/utility/selection_policy.py: make_policy(args.select) returns LATENCY or POWER, and tuners get it as self.policy / self.policy_for(args). It wraps the existing helpers: rank, beats (with an optional latency noise margin for the MoE flat check), best(candidates, latency_choice=...) for tuners with their own latency rule (the MXFP4 finalist gate), sort_frame, pair_energy, accurate_count, plus output helpers (attach_columns, result_metrics, watt_text, profile_file). rank_moe_pairs, the grouped, MXFP4, and staged loops, K5, IQ2R, a16w16, gradlib, and base_tuner all call the policy. The staged MXFP4 search now ranks and measures through it too; before this it compared us only.
  2. Measurement in one place. Tuners hand policy.measure(call) / policy.record(target, key, call) the kernel launch. mp_tuner calls policy.measure_timed(func, args, kwargs). The latency policy never touches the sensor and adds no watts column.
  3. AITER_SELECT is runtime-only. The tuner no longer writes it, and --select no longer defaults from it. mp_tuner takes policy= or reads the one TunerCommon.run activates (use_policy), so the ~17 existing mp_tuner callers needed no change. jit/core and power_configs are the only readers. This also stops AITER_SELECT=power from leaking into the --compare / --run_config benchmarks during a tune.

test_tuners_do_not_decide_select_mode_themselves enforces this across base_tuner, mp_tuner, gradlib, and every csrc/**/*_tune.py: none of them may call power_selection_enabled, select_power, the sensor functions, the raw rank helpers, or AITER_SELECT. A few tuners still read policy.measures_power only to decide whether to print watts columns in their log output.

Tuner unit tests pass locally (50). test_tune_pipeline and test_run_config need Ninja for the JIT build, which this host lacks, so I'm relying on CI for those.

@aiter-bot

aiter-bot commented Oct 29, 2026 •

Copy link
Copy Markdown

READY FOR HUMAN REVIEW

PR Self-Review

Revision bbc0ae5a8954 · mode shadow

Author self-review accepted. Readiness report:

Author answers

  • PASS — current-SHA answers are present and structurally specific
Show the author's self-review answers

Behavior and intent

Tuners rank candidates by latency only, so the tuned tables always pick the fastest kernel even when a slightly slower one uses much less package energy. This PR adds an opt-in `--select power` mode that ranks accuracy-qualified candidates by `us × package watts` and stores those rows separately (`aiter/configs/power_configs/`), loaded only with `AITER_SELECT=power`. Default behavior is unchanged: `--select latency` keeps the lowest `us` and writes no power columns. Supported tuners: the shared GEMM tuners (`GemmCommonTuner`: a16w16 incl. hipBLASLt/asm/FlyDSL/Opus/Triton/torch, a8w8, blockscale, gradlib), the fused-MoE CK two-stage, grouped, and FlyDSL MXFP4 paths, the GDN K5 BV sweep, and IQ2R. Conv3d, a6w6, and the mixed-MXFP tuners (a6w4/a4w6) reject `--select power`, because their gain guard keeps a kernel only for speed.

Validation and unverified scope

Hardware: 8× MI355X (gfx950), ROCm 7.2.4, power cap 1400 W, clocks auto (SCLK 500-2400 MHz), GPUs otherwise idle.
- Unit: `python -m unittest op_tests.tuning_tests.test_fast_tuning` — 54 tests pass, including a real-sensor test (8192³ bf16 GEMM must read > 1.5× idle and ≥ 0.85× a back-to-back loop; skips without `power1_input`).
- Other tuning suites pass: `test_mp_tuner_logic`, `test_compare_logic`, `test_config_shape_collision`, `test_fhmoe_tuner`, `test_a6w6_tuning`, `test_mixed_mxfp_tuning`. `test_tuner_infra` has 2 failures (`a_preshuffle` blank-cell) that also occur on main (#6709).
- End to end, latency and power tunes on GPUs 0-7, low/mid/high shapes from the shipped untuned CSVs: a16w16 (all backends + hipBLASLt, 4 shapes), gradlib (4), a8w8 (3), a8w8 blockscale (3), CK MoE ASM FP8 (3), MXFP4 staged (3) and fast scan (2), IQ2R (3), and a16w16 with `--mp 2`. On every shape the latency CSV has no power columns, every power row has finite watts and the conditions, and the power pick is the minimum-energy measured candidate. Examples: a8w8 16384×8192×1024 is 23% slower and 9% lower energy; blockscale 1024×1536×7168 is 24% slower and 16% lower energy; MXFP4 32 tokens is 4% slower and 18% lower energy.
- MoE profile schema (FP8 per_1x128, ASM, 3 shapes): the latency `-o2` profile has main's 29 columns with no power columns, and `profile_fmoe.csv` has none; the power run keeps the energy-ranked profile and stamps all rows.
- Power sample vs CUDA-graph replay (no launch gaps): `torch.mm` within ~10%; split-K FlyDSL still reads 12-28% low after the launch-gap correction, though the energy ranking on that shape matches the graph-replay ranking.
- Unverified: grouped MoE (gfx1250 only) is covered by unit tests only; K5 has no shipped untuned shapes; gfx942 not run; kernels under ~5 µs are left unmeasured by design; the power score is an isolated replay, so a winner can change in a model, at another power cap, or next to other kernels; end-to-end model perf/W (MLPerf) is not part of this PR. Graph-replay power sampling is planned as a follow-up.

PR-specific question

No kernel dispatch path changes. Which kernel runs is still chosen by the tuned CSV for the arch/dtype/shape keys, as before. A power row can reach dispatch only through `AITER_SELECT=power` (overlay of `aiter/configs/power_configs/`) or an explicit `AITER_CONFIG_*` path. Without `AITER_SELECT=power`, `power_configs/` is ignored. A power row is written only for a candidate that passed the same `errRatio` accuracy gate as latency mode and has finite measured watts; a batch with no finite watts stays a latency result. Merging a power row and a latency row for one shape, or power rows tuned under different perf level, cap, or SCLK, raises instead of auto-resolving. Tuners that cannot rank by energy reject `--select power` (`SUPPORTS_POWER_SELECT = False`).

Changed op tests

`op_tests/tuning_tests/test_fast_tuning.py` — the only changed op test. Its new cases cover power selection, and none duplicate an existing case:
- Ranking and policy: energy vs latency order, unmeasured rows ranked last, MoE stage-energy pairs, `candidate_beats`, `SelectionPolicy` decisions, `mp_tuner` using the activated policy, a16w16 passing its policy to hipBLASLt.
- Measurement: the timed kernel is replayed without `run_perftest` options; bounded in-flight launches; launch-gap correction; real-sensor large-GEMM check.
- Output and loading: only measured rows are kept and stamped, power output redirected to `power_configs/`, latency frames and the latency MoE profile without power columns, power/latency merge conflicts, the `power_configs/` overlay (unstamped rows and files without `select` ignored).
- Build safety: `jit.core` loaded as in `setup.py` (with and without duplicate shapes) does not import the `aiter` package.
- Guard: no tuner file decides the select mode, reads power, or uses `AITER_SELECT` itself.

PR state

  • PASS — pull request is marked ready for review

Copilot findings

  • PASS — no active Copilot finding lacks a substantive author reply
  • ADVISORY — Copilot did not review the current head SHA

Code and op-test hygiene

  • PASS — no committed artifact or conflict marker detected

Required CI

  • PASS — Black: passed
  • PASS — Ruff: passed
  • PASS — Repository dependency: passed
  • PASS — Aiter Test Gate: passed

Shadow mode reports what would block but keeps the status successful. Set AITER_PR_SELF_REVIEW_MODE=enforce only after reviewing real false positives.

@yzhou103

yzhou103 commented Oct 30, 2026 •

Copy link
Copy Markdown
Contributor

Thanks for the update. I checked that latency mode is unchanged versus main (849f711), with AITER_SELECT unset:

  • All 27 runtime config files are byte-identical.
  • post_process and the merge give identical output on fixed synthetic results, including ties and invalid times.
  • Real a16w16 and fused-MoE latency tunes pick the same kernels, with the same tuned-CSV schema.
    I also ran gemm_a8w8_mxfp8_bpreshuffle with --select power. Its own result_to_df override still gets watts, tflops_per_w, select=power and the power-state columns, so centralizing these in post_process and run() works.

Please fix in this PR:

  1. MoE profile schema changes in latency mode. The -o2 profile file now gets rank, watts1, watts2, energy, watts and tflops_per_w, and profile_fmoe.csv gets watts. All of them are NaN in latency mode, and the -o2 column and row order also differ from main. The tuned CSV is unchanged, but anything that reads the profile files sees a different schema. Please add these columns only when select=power, and keep the latency output the same as main.

  2. Unmeasured rows are written to power_configs. In a power tune, shapes whose watts could not be measured fall back to a latency rank. _stamp_power_rows leaves select and the power-state columns blank on those rows, but still returns them, so they are written to aiter/configs/power_configs/. When no row in a batch is measured, the whole batch is written ("leaving this batch as a latency result"). At runtime, _overlay_power_tables replaces the shipped latency row with any power_configs row regardless of select, so with AITER_SELECT=power these latency-ranked rows override the shipped (or model_configs) rows.

    Suggested fix:

    • In _stamp_power_rows, keep only measured rows and log a warning with the skipped shape keys. If nothing is measured, return an empty frame. Skipped shapes then show up as untuned and are retried on the next power tune.
    • In _overlay_power_tables, keep only rows where row_select_mode(select) == "power" before deduplicating, and ignore a power file that has no select column. This protects against older or hand-copied tables.

Everything else looks good to me.

@yzhou103 yzhou103 closed this Oct 30, 2026
@yzhou103 yzhou103 reopened this Oct 30, 2026
@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:gfx1250-ffm-triton Run the five-shard gfx1250 FFM Triton test suite
ci:triton-300x Run an additional Triton test job on MI300X in PRs (added automatically when gfx942 configs change); main branch always runs both MI35X and MI300X
ci:triton-355 Run the full Triton test suite on MI35X, not only the tests the change affects
ci:windows Windows JIT op smoke test on the gfx1201 ROCm runner
ci:multi-gpu Aiter multi-GPU tests on the 8-GPU runner
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: DeepSeek-V4-Pro, MiniMax-M3-MXFP4
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 6970 --add-label <label>

One backend per PR:
A PR changes one kernel backend: [Triton/Gluon] (Triton and Gluon count as one), [HIP], [ASM], [CK], [OPUS] or [FlyDSL]. If the title ends up with two backend tags, split the PR -- as stacked pull requests when one part cannot merge without the other.

PR title tags & labels:
Component tags ([Triton/Gluon], [HIP], [CK], [ASM], ...) are added to the PR title and as PR labels automatically from the changed files and re-synced on every push — change-type tags like [fix]/[Perf], op tags like [MLA], and human labels (ci:*) are left untouched. Add the no-auto-title label to stop the title rewrites; labels stay in sync either way.

@rpoornac

Copy link
Copy Markdown
Contributor Author

Thanks @yzhou103 for checking latency mode against main. Both are fixed in 6fae460.

1. MoE profile schema in latency mode. Stage rows carry watts only with --select power, and rank_moe_pairs returns the frame unchanged in latency mode. The -o2 file starts from self.columns as on main, and the pairs are written in main's order. The latency winner is still the lowest-us row (same as main's idxmin). FP8 per_1x128 MoE tune on MI355X (ASM, 1/64/4096 tokens): the latency -o2 profile went from 35 to main's 29 columns, with no rank, watts1, watts2, energy, watts, or tflops_per_w, and profile_fmoe.csv has no watts. The power run still writes the energy-ranked profile and stamps every row.

2. Unmeasured rows in power_configs. I took your suggestion:

  • _stamp_power_rows keeps only measured rows and logs the skipped shape keys. If nothing is measured, it returns zero rows (columns kept, because K5's result_to_csv needs them). Those shapes are reported as untuned and retried by the next power tune.
  • _overlay_power_tables uses only rows where row_select_mode(select) == "power", so a file without a select column contributes nothing, and it logs how many rows it ignored. The filter runs after the duplicate-shape check, because that check can rewrite source files row by row, and filtering first would delete the ignored rows from the user's file.

New tests: latency rank_moe_pairs keeps its columns and order; mixed and all-unmeasured batches drop the unmeasured rows; the overlay ignores unstamped rows and files without select. test_fast_tuning (54), test_config_shape_collision, test_fhmoe_tuner, test_mp_tuner_logic, and test_compare_logic pass locally.

rpoornac and others added 4 commits November 1, 2026 11:06
Add --select power so shared tuners keep the lowest runtime times PPT watts. The accuracy gate is unchanged, and latency remains the default.

Co-authored-by: Cursor <cursoragent@cursor.com>
Format the tuner edits and replace the NaN self-comparisons and bare exception handlers that the pre-check jobs reject.

Co-authored-by: Cursor <cursoragent@cursor.com>
Sort power candidates by energy and then latency so an unread sensor cannot drop the fastest kernel. A single AITER_POWER_PATH is rejected when more than one tuner process is running.

Co-authored-by: Cursor <cursoragent@cursor.com>
Power tunes of shipped tables now write to aiter/configs/power_configs/, which loads only with AITER_SELECT=power and replaces only the shapes it covers. Each power row records perf level, power cap, and SCLK range, and the config merge refuses to auto-resolve power rows against latency rows or rows tuned under different conditions. Power mode is marked experimental.

Co-authored-by: Cursor <cursoragent@cursor.com>
rpoornac and others added 9 commits November 1, 2026 11:06
Ranking, package-power measurement, and power_configs loading each have one owner. Latency CSV columns stay unchanged unless --select power.

Co-authored-by: Cursor <cursoragent@cursor.com>
The replay was passing run_perftest keywords into the kernel, so watts came back NaN and the row was still stamped select=power. A batch with no finite watts stays a latency result, and --select power is rejected on tuners that do not record it.

Co-authored-by: Cursor <cursoragent@cursor.com>
Main now rejects an MXFP4 row before the scan unless dtype and quant fields match the port. The energy-ranking fixture was missing those fields, so it never reached the retime.

Co-authored-by: Cursor <cursoragent@cursor.com>
CI skipped the GPU tests because Black rejected these three files.

Co-authored-by: Cursor <cursoragent@cursor.com>
setup.py loads jit.core before aiter exists. Importing the power overlay from that path ran aiter/__init__.py while module_aiter_core was still being registered.

Co-authored-by: Cursor <cursoragent@cursor.com>
A tuner builds a SelectionPolicy from --select and asks it for the
candidate rank, frame order, pair energy, accurate-pass count, and the
power sample. Tuners no longer branch on the select mode or write
AITER_SELECT; that variable only picks the tables jit/core loads.
mp_tuner reads the policy TunerCommon.run activates.

package_power.measure_kernel is the one power read. mp_tuner hands it
the kernel launch with the run_perftest options removed, so the window
replays the kernel back to back. Replaying run_perftest itself left the
GPU idle for most of the window.

The duplicate-shape dedup in jit/core sorts by lowest us without
importing the aiter package when setup.py loads it as jit.core.

Co-authored-by: Cursor <cursoragent@cursor.com>
- TunerCommon.run wrapped its body in _run, which IQ2R already defines
  as its kernel launcher, so every IQ2R tune failed. Use _run_selected.
- The a16w16 tuner built gradlib's Gemm without the policy, so hipBLASLt
  was never power-measured and always ranked last. Pass the policy, and
  default Gemm and the K5 sweep to the active policy.
- The power window queued launches without bound. IQ2R graph replays
  crash inside the HIP runtime when hundreds are queued. Keep a few
  events in flight, one per ~200 us of work, so the GPU never drains.
- IQ2R times a captured CUDA graph; sample power from that graph too.
- Staged MXFP4 logs each pair's watts like the other loops.

Co-authored-by: Cursor <cursoragent@cursor.com>
A kernel whose Python launch takes longer than its GPU time leaves the
GPU idle between launches, so the window read close to idle and the
kernel looked efficient (split-K FlyDSL GEMMs at 256x4096x4096 read
~370 W against ~990 W replayed from a CUDA graph, and won power mode).

Every measurement site now passes the call's timed us. When the loop
never waited on the GPU, the window's busy fraction is launches * us /
window, and the reading is scaled to active power over the between-launch
baseline (a one-element kernel loop, ~300 W on MI355X). A window that
waited on the GPU kept it saturated and stands as read. Under 20% busy
the candidate stays unmeasured.

Co-authored-by: Cursor <cursoragent@cursor.com>
- MoE latency tunes write the -o2 profile and profile_fmoe.csv with
  main's columns and row order. Stage watts, rank, energy, and
  tflops_per_w appear only with --select power.
- A power tune drops rows whose watts were not measured instead of
  writing latency picks to power_configs/. Those shapes stay untuned
  and are retried by the next power tune.
- The AITER_SELECT=power overlay uses only rows stamped select=power,
  so an older or hand-copied table cannot replace shipped rows.

Co-authored-by: Cursor <cursoragent@cursor.com>
@rpoornac
rpoornac force-pushed the power-kernel-select branch from 93678ee to c83abcd Compare November 1, 2026 17:07

@yzhou103 yzhou103 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.

LGTM

@yzhou103
yzhou103 merged commit b50a0d2 into ROCm:main Nov 3, 2026
57 checks passed
amd-bartgips added a commit that referenced this pull request Nov 3, 2026
…status callers

#6970 attaches package watts to each result as TuneResult.watts. A
return_status=True result was a plain 5-tuple, so under --select power its
watts read as NaN and the tune wrote no rows for that caller.
TuneStatusResult still unpacks as (info, us, err_ratio, status, detail).
amd-bartgips added a commit that referenced this pull request Nov 4, 2026
…status callers

#6970 attaches package watts to each result as TuneResult.watts. A
return_status=True result was a plain 5-tuple, so under --select power its
watts read as NaN and the tune wrote no rows for that caller.
TuneStatusResult still unpacks as (info, us, err_ratio, status, detail).
vgokhale added a commit that referenced this pull request Nov 5, 2026
* CI: auto-update split test times (#7290)

Co-authored-by: gyohuangxin <42127654+gyohuangxin@users.noreply.github.com>

* [HIP] [FlyDSL] perf(mega_moe): fuse the gfx1250 compact dispatch and GEMM1 into one stage-1 kernel (#6564)

* [ASM] [HIP] [OPUS] Add gfx950 MLA v4.1 a8w4 prefill kernels for an NVFP4 prefix KV (#7124)

* Add a gfx950 MLA v4.1 a8w4 prefill kernel for an NVFP4 prefix KV

* Add a gfx950 MLA v4.1 a8w4 16mx4_64nx1 prefill kernel and run the NVFP4 prefix in f16

* Run the dsa_v41 a8w4 NVFP4 prefix in bf16 to keep MXFP8 Q exact

* [Perf] Avoid torch Stream construction across grouped MoE launches (#7219)

* Avoid constructing a torch Stream for default grouped GEMM launches

* Reuse raw stream arguments through MoE route, quant and reduction wrappers

* [Bugfix] MLA decode: merge only the splits gfx950 bf16 gqa32 stage 1 writes (#7182)

On gfx950, mla_a16w16_qh16_m16x4_n32x1_coex0_mask1 (bf16, gqa 32,
non-persistent) deals KV to splits round-robin in 32-token tiles and
exits without writing a split that gets no tile. mla_decode_fwd used
mgc = 16 for nhead 32, so stage 2 merged min(splits, cdiv(kv, 16))
splits and read up to twice as many as stage 1 wrote, from torch.empty
buffers. Rows with less KV than splits * 32 came back wrong, or NaN
when that memory held NaN. Use mgc = 32 for this kernel, as the fp8
qh64 kernel already does with mgc = 64. gfx942's gqa32 kernel uses
16-token tiles and keeps mgc = 16.

The new op test runs that case with NaN-filled split buffers.

Signed-off-by: Batzorig Zorigoo <baz@fractalyze.io>

* [Config] [gfx1250] Tune MiniMax-M3 TP4 BF16, indexer-CP QKV and MoE token 1024 (#7209)

* [Triton/Gluon] Window-aware split-KV for sliding-window verify on gfx950 (#6889)

* [Triton/Gluon] Split window-trimmed sliding-window verify on gfx950

On gfx950 a sliding-window call with q > 1 never takes the split grid
(_gfx950_split_grid returns False, and kv_split["SW"] caps MAX_SEGMENTS
at 1), so a spec-decode verify on a sliding layer runs num_kv_heads
programs per sequence over the whole window. Splitting it only pays once
the KV range is cut to the window: the split carves seq_len, so over a
full context the window lands in one or two segments and the rest are
empty (31.9-37.6 us vs 29 us unsplit when forced).

Split only when the KV range is bounded by the window. Row i of a q-row
sequence of length L reads keys from L - q + i - (SW - 1) on, so a block
table trimmed to whole pages holds at most SW + q - 1 + block_size - 1
keys. The new SWB config axis is sliding_window_bounds_kv(): q > 1 and
max_seqlen_k <= sliding_window + max_seqlen_q + block_size - 2. A batch
whose contexts are all that short is bounded as well, trimmed or not, and
takes the same path. _gfx950_split_grid lets verify-sized SWB calls
(max_seqlen_q <= 16 and max_seqlen_q * num_queries_per_kv <= 128)
through and then decides like any other batch (NUM_SEGMENTS > 1).
Longer queries (prefill chunks) keep the 2d grid: with q rows per
sequence in the reduce, MHA q = 128 at window 256-512 ran up to 1.33x
slower split.

The gfx950 gluon kv_split table gains an "SW.SWB" entry, the SW entry
without its segment cap, plus:
- SPLIT_FIT_BUDGET (new): the largest power-of-two split that keeps the
  launch within SEGMENTS_PER_CU (1) programs per CU, sized by the real
  one-block-per-sequence query grid (_gfx950_split_params). The default
  ceil-rounded share splits 2 ways up to B = 51 and loses 30-40% from
  B = 36 (39.7-43.0 us vs 30.4-30.8 us on the 2d grid).
- SPLIT_KV_WINDOW (new): count the KV tiles over the sliding_window +
  max_seqlen_q - 1 keys the rows can see, not the page slack of a trimmed
  table, so SPLIT_MIN_TILES 4 keeps window 128 unsplit at block 128 too
  (64 q / 8 kv heads, q = 4: 8.2 us split vs 6.7 us 2d).
- SPLIT_SEQS_PER_TILE 1.5 (new): no split once num_seqs exceeds 1.5 x the
  KV tiles. The triton reduce_segments pass that every gfx950 split pays
  costs more with the batch (7.6 us at B = 32, 2 segments), while the 2d
  kernel's time follows the KV length, so at windows 256-1024 a split at
  B = 12-32 ran up to 1.87x slower than the 2d grid.

Kernel: a split whose tiles all fall before the window start stores an
empty partial, and split 0 is the only one that counts the attention
sink, so a window starting past tile 0 (block size > 64, or a shorter
row in an untrimmed bounded batch) dropped the sink from the softmax.
The empty-partial path of split 0 now stores M = sink, L = 1.

Unchanged dispatch: SW calls with a context past the bound, SW decode
(q == 1 keeps the cap), q > 16, window 128, full attention, gfx1250 and
the triton backend. Of 328,800 no-launch shapes (window 128-32768, q
1-128, B 1-128, 5 head layouts, D 64-256, bf16/fp8, block 16-256,
ctx bound -300..+5000) only gluon SW calls with 1 < q <= 16 inside the
bound change grid (22,392).

MI355X, 16 q / 2 kv heads, head 64, bf16, block 16, window 2048, q = 8,
ctx 32768 (us per call, CUDA-graph replay, natural dispatch):
  B           1     2     4     8    16    32    48    64
  full      29.0  29.0  29.1  29.5  30.0  30.5  30.6  30.9
  trimmed    6.7   7.2   8.6  10.9  15.8  25.1  30.6  30.8
  segments    64    32    16     8     4     2    2d    2d
Full-attention floor at 2056 keys: 6.6 / 6.8 / 8.1 / 9.4 / 14.4 / 21.9.
The natural pick matches the fastest forced count (1-128) at every B.
Shorter windows, trimmed, same heads (base 2d -> split, us):
  W 512:  B4 9.78 -> 6.79, B8 9.79 -> 7.61, B12 9.82 -> 9.19, B16+ 2d
  W 1024: B4 16.30 -> 7.66, B8 16.35 -> 8.61, B16 16.71 -> 12.27, B32 2d
  W 256:  B4 6.54 -> 6.30, B8+ 2d

Known cost: a caller that captures a graph at max_seqlen_k = window bound
and replays it for batches whose real contexts are a few hundred keys
pays the split + reduce where the 2d kernel was cheaper (B >= 12, ctx <=
1000). A cheaper gfx950 reduce would remove that.

Test: test_triton_unified_attn_sw_verify_split (full and trimmed tables
vs reference with a dominant sink, exact segment counts: trimmed verify
splits, full context stays 2d, SW decode stays 1 segment, q = 17 and
q = 128 prefill chunks stay 2d, block 128 and an untrimmed bounded batch
leave split 0 empty) and test_unified_attn_sw_kv_split_config
(CPU-only kv_split picks at the bound edge, untrimmed, decode, window 128,
and the batch sizes where the split stops).

Signed-off-by: Oxana Korzh <okorzh@amd.com>
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

* [Triton/Gluon] Drop SPLIT_KV_WINDOW; clamp every window-bounded split

Review follow-up. The SPLIT_KV_WINDOW knob was set only in the SW.SWB
entry, so it configured nothing. compute_segment_params now counts KV
tiles over the sliding_window + max_seqlen_q - 1 keys the rows can see
whenever sliding_window_bounds_kv() holds, with no config switch.

The clamp is gated on the bound, not on sliding_window > 0. The split
carves the whole seq_len, so on an untrimmed long context a window-sized
count leaves most segments idle. No unbounded sliding-window call reaches
a split lookup today, so the two gates pick the same grid.

CPU dispatch sweep vs the previous commit, 2,048,256 shapes per arch on
gfx950 (gluon and triton), gfx1250, gfx1100 and gfx1201: no 2d/3d,
NUM_SEGMENTS or TILE_SIZE change. 1,184 gfx950 gluon shapes change only
the triton reduce num_warps (head 256, block >= 64, window + q - 1 of
about 257-512), since the reduce lookup now sees the same clamped length.
test_unified_attn_sw_kv_split_config: 11 passed on CPU; removing the clamp
fails the window 128, block 128 case.

Signed-off-by: Oxana Korzh <okorzh@amd.com>
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

* [Triton/Gluon] Key the SWA split on window size: numeric SW axis

Review follow-up. The SW axis was a bool, so a table could not target a
window size. SW is now "opt_num": the keys a row can see,
min(sliding_window, max_seqlen_k), floored at 1, and 0 with no window.
SW_EQ/LEQ/GEQ_<n> keys pick by size, a bare SW matches any window, and
no window matches only "any". Tables that only use a bare SW key (gfx950
gluon and triton, gfx1250 gluon, gfx1201 triton) resolve as before.
_index rejects a table that mixes SW_LEQ and SW_GEQ keys, because every
matching LEQ key would outrank every GEQ key. It is cached on the key
names, so the check runs once per table.

The gfx950 gluon kv_split entry SW.SWB becomes SW_GEQ_256.SWB. Bare SW,
with MAX_SEGMENTS 1, stays the fallback for every other window and for SW
decode. 256 keys is 4 tiles of 64, the smallest window measured to win.
Window 128 now falls back through the threshold, not SPLIT_MIN_TILES.
SWB stays: an untrimmed long context also gives SW = window, and
splitting that table was measured slower.

CPU dispatch sweep vs the previous commit, 3,692,160 shapes per arch
(grid, NUM_SEGMENTS, reduce config, and attn_2d/attn_3d configs for both
backends) on gfx950, gfx1250, gfx1100 and gfx1201: the only change is
6,300 gfx950 gluon shapes at window 255 (q 4-16, just under the
threshold), split -> 2d. Windows above 2048 still split; their perf is
not measured.

Tests: test_unified_attn_sw_kv_split_config gains window 255 vs 256 at
the threshold and window 256 at block 128 (the clamp case).
test_unified_attn_sw_axis_lookup covers the candidate order, the axis
value and its floor, and the LEQ/GEQ guard. 15 passed on CPU. Each of
these mutants fails at least one test: threshold, clamp, guard, bare-SW
fallback, the no-window check, min(), and the floor.

Signed-off-by: Oxana Korzh <okorzh@amd.com>
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

* [Triton/Gluon] Test the SWA split without sinks on long pages

Review follow-up (Copilot). Without sinks, a sliding-window split over a
bounded table whose page is much longer than a tile can leave 32 or more
empty leading segments. One example is head 256, q 8, KV 10246, window
2048, block 8192, 64 segments. Before #6888 that made the first
SEGMENT_BLOCK chunk of reduce_segments all -inf, and the merge returned
NaN. #6888's two-pass merge takes the global max first, so those chunks
now weigh 0. Under TRITON_INTERPRET=1, main's reduce_segments on the
partials both shapes leave gives no NaN and matches the single-chunk
reduce within 1.2e-7.

test_triton_unified_attn_sw_split_no_sinks runs both of Copilot's shapes
on gfx950 with head 256 and no sinks, against the reference.

Signed-off-by: Oxana Korzh <okorzh@amd.com>
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

* [Triton/Gluon] SWA split tests: check accuracy only

Review follow-up. Tests here target accuracy, not configs or routing.

- Remove test_unified_attn_sw_kv_split_config and
  test_unified_attn_sw_axis_lookup. They pinned kv_split picks and lookup
  order, which couples the tests to the tuned configs and heuristics.
- test_triton_unified_attn_sw_verify_split and
  test_triton_unified_attn_sw_split_no_sinks no longer record or assert
  the grid. Both still compare full and trimmed tables, or the no-sink
  long-page shapes, against the reference.

Signed-off-by: Oxana Korzh <okorzh@amd.com>
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

* [Triton/Gluon] Window-aware split-KV for sliding-window attention

With a sliding window, every split-KV launch carved [0, seq_len) into
NUM_SEGMENTS tile ranges, so over an untrimmed long context only the few
segments that overlap the window did work. That is why this PR so far
split only tables bounded by the window.

Carve only the window instead. For SLIDING_WINDOW > 0 each sequence
splits the tiles [tile_lo, cdiv(seq_len, TILE_SIZE)), where tile_lo is
the tile of row 0's first visible key, max(context_len - SLIDING_WINDOW
+ 1, 0) // TILE_SIZE. It is per sequence, so every query block and the
reduce agree. The split launch already puts a sequence's spec-decode rows
in one query block when q * num_queries_per_kv <= 128. The window's tiles
are now what gets split across the segment workgroups.

- The gfx950 gluon split, kernel_unified_attention_3d and reduce_segments
  carve this way. The triton pair shares split_window_tiles, and
  reduce_segments gets SLIDING_WINDOW from the launch. The gfx1250 3d
  wrapper asserts no window: its split still carves seq_len, and
  use_2d_kernel never sends a window there. With no window, all three
  kernels compile to the same gfx950 instructions as before.
- Dispatch: an untrimmed table now splits like a trimmed one.
  _gfx950_split_grid drops the bound check and keeps the q <= 16 and
  q * num_queries_per_kv <= 128 gates. The kv_len clamp applies to every
  sliding-window call. The SWB axis and sliding_window_bounds_kv are gone,
  and the split entry is SW_GEQ_256.Q_GEQ_2, so sliding-window decode
  stays on the capped bare SW entry. CPU dispatch sweep vs the previous
  commit, 3,692,160 shapes per arch: on gfx950 the only change is 75,220
  gluon sliding-window calls with 1 < q <= 16 past the old bound going
  2d -> split. Each gets exactly the dispatch the bounded table already
  had. gfx1250 and gfx1100 do not change.

The existing accuracy tests now run the full-table rows through the
window carve: test_triton_unified_attn_sw_verify_split and
test_triton_unified_attn_sw_split_no_sinks.

Signed-off-by: Oxana Korzh <okorzh@amd.com>
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

---------

Signed-off-by: Oxana Korzh <okorzh@amd.com>
Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
Co-authored-by: Satya Nikhil Kodukula <nikhil.kodukula@gmail.com>

* [FlyDSL] gfx1201 (RDNA4) FlyDSL BF16 attention optimizations & FP8 attention implementation (#4188)

* feat(flydsl): add per-tensor fp8 gfx1201 flash-attention kernel

* feat(flydsl): generalize the gfx1201 bf16 flash-attention kernel and wire fp8 dispatch

* feat(flydsl): fully-FlyDSL fp8 per-tensor quant+rotation for gfx1201

* refactor(flydsl): share and harden gfx1201 attention

* test(flydsl): cover gfx1201 bf16 and fp8 attention

* bench(flydsl): add gfx1201 fmha benchmark

* merge(main): resolve gfx1201 FlyDSL attention conflict

* fix(flydsl): validate attention softmax scale

* fix(flydsl): address gfx1201 attention review

* refactor(flydsl): centralize gfx1201 attention build parameters

* test(flydsl): execute gfx1201 pytest suite in op-test shards

* refactor(flydsl): share gfx1201 attention helpers and localize FP8 intrinsics

* refactor(flydsl): clean up gfx1201 kernels

* refactor(flydsl): organize gfx1201 FMHA modules

* refactor(flydsl): localize gfx1201 stream readiness

* refactor(flydsl): use typed pointers in fp8 quant

* refactor(flydsl): use public gfx1201 kernel attributes

* refactor(flydsl): export gfx1201 fp8 quant directly

* test(flydsl): keep gfx1201 FMHA public regressions

* feat(flydsl): add native gfx1201 1 FP8 quantization for more shapes

* flydsl: simplify gfx1201 FMHA stream/quant integration

* op_tests: modernize gfx1201 FMHA validation and benchmark

* flydsl: harden gfx1201 FMHA stream validation and multi-device tests

* flydsl: remove empty gfx1201 FMHA package init

* flydsl: migrate gfx1201 FP8 flash-attn to FlyDSL buffer-view/copy API

* flydsl: use static layouts and coordinate helpers for gfx1201 FP8 KV loads

* flydsl: validate gfx1201 FP8 buffer indices against signed Int32 limits

* flydsl: replace raw arith intrinsics with typed FlyDSL scalar/vector ops in gfx1201 FP8 FMHA

* refactor(flydsl): use typed exp2 in gfx1201 BF16 attention

* test(flydsl): tighten gfx1201 FMHA validation

* refactor: modernize gfx1201 FlyDSL FP8 paths

* test: consolidate FlyDSL gfx1201 op tests

* fix: support wide FlyDSL FP8 quantizer loads

* fix: validate FlyDSL quant offsets and stream coverage

* fix(flydsl): harden gfx1201 FMHA correctness

* fix(flydsl): align FP8 softmax normalization to prevent long-sequence normalization bias

---------

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

* [Triton/Gluon] Enable and tune pipelined unshuffled MLA decode on gfx1250 (#6545)

* Allow the pipelined version of the kernel to be used

* Add tuning config

* More tuning changes

* Pad fp8 rows

* Update reduce waves_per_eu

* Fit 64 heads into two waves

* Add unit test

* Formatting

* Attend to Copilot review comments

* Test fixes

* Update tuning params

* Attend to Copilot review comments

* Fix formatting issues

* Add extra guards

* Add params to desc

* Fix arch name for test

* [Triton/Gluon] [Kernel] Add dual-layout H16 MXFP4 quantization (#6458)

* [Triton/Gluon] [Kernel] Add dual-layout H16 MXFP4 quantization

Fuse row-wise MXFP4 quantization with signed-H16 quantization of the
transposed input. Support independent stochastic rounding and GEMM
layouts on gfx950, with tests, benchmarks and a tuning configuration.

Signed-off-by: DaiXindi-AMD <xdai@amd.com>

* [Misc] Rename fused MXFP4 tuning configuration directory

Signed-off-by: DaiXindi-AMD <xdai@amd.com>

* [Bugfix] Address fused transpose Hadamard MXFP4 review feedback

Use the specific fused API and config key, relocate the four unchanged device helpers to quant.py, and document sign validation and physical scale shapes. Fix module-object monkeypatching after the API/module rename.

Validated all eight changed Python files with Ruff 0.16.0, Black 26.3.0 and 26.10.1, compile and diff checks. Kernel and helper AST equivalence plus isolated CPU helper checks passed. Full package import and GPU tests were not rerun: prior import/collection is blocked by missing Ninja, and no GPU allocation is authorized.

Signed-off-by: DaiXindi-AMD <xdai@amd.com>

* [Docs] Use GPU graph capture terminology for H16 signs

Signed-off-by: DaiXindi-AMD <xdai@amd.com>

* [Misc] Keep MXFP4 layout constants in the host utility

Move the five unchanged layout constants to utils/mxfp4.py and remove the constants-only utils/_triton/mxfp4.py module. The public wrapper still passes the same values as constexpr arguments; device kernel and helper code remain unchanged.

Signed-off-by: DaiXindi-AMD <xdai@amd.com>

---------

Signed-off-by: DaiXindi-AMD <xdai@amd.com>

* [Triton/Gluon] [Perf] fused_reduce_act_mul_and_mxfp4_quant: wider launch tile, add round_to_input_dtype and transpose_scale (#6024)

* [Triton/Gluon] [Perf] fused_reduce_act_mul_and_mxfp4_quant: wider launch tile, add round_to_input_dtype and transpose_scale

For N_half > 1024 (no shuffle, no scale padding, no split-K) launch one warp
per row over 256 (M <= 64) or 512 columns instead of a 1x128 tile on 4 warps.
On MI355X thats 1.02-3.77x faster over N_half 2048..28672 and M 1..8192
(9.85 -> 4.29 us at M=64, N_half=14336), output bytes unchanged.

round_to_input_dtype rounds act(x) * x_mul to the input dtype before the
quant, so the result is bitwise equal to act_mul in that dtype followed by
dynamic_mxfp4_quant. transpose_scale (also added to fused_rms_mxfp4_quant)
returns the scales column-major, the layout dynamic_mxfp4_quant returns.
Both default to False.

Signed-off-by: adlashab <adlashab@amd.com>

* Fix a typo in the act_mul launch tile comment

Signed-off-by: adlashab <adlashab@amd.com>

* [Triton/Gluon] fused_reduce_act_mul_and_mxfp4_quant: honor scaling_mode="ceil"

scaling_mode was accepted but the act_mul kernel always used the "even"
scale. Pass it through to _mxfp4_quant_op, which already has "ceil" as
SCALING_MODE 1: 2^ceil(log2(amax / 6)), the RoundUp scale
per_1x32_f4_quant_hip uses by default. With round_to_input_dtype the output
is then the same fp4 bytes and the same scales as silu_mul followed by
per_1x32_f4_quant_hip, in the same scale layout (plain, padded or shuffled),
so it can stand in for that quant in front of gemm_afp4wfp4_preshuffle /
gemm_a4w4. That is what vLLM's MXFP4 linears on the ASM layout run. One
difference: a 32-block thats all zeros gets scale byte 1 here and 92 from the
HIP kernel. All its fp4 values are 0, so the GEMM output is the same.

Also let the wider launch tile run with scale_shuffle_padding when shuffle is
off. The padding only changes the allocation, the scale store is the same.

New test against per_1x32_f4_quant_hip for plain, padded and shuffled scales.

Signed-off-by: adlashab <adlashab@amd.com>

* [Triton/Gluon] [Test] Check fused act_mul + MXFP4 quant against a torch reference

Add a plain torch reference for the per-1x32 MXFP4 quant and compare
fused_reduce_act_mul_and_mxfp4_quant against it bit for bit. The reference
uses the shared CPU helpers from aiter.utility.fp4_utils: the e8m0 scale from
f32_to_mx_e8m0_scale (Even for "even", RoundUp for "ceil"), then x / scale
rounded to nearest even on the E2M1 grid with f32_to_mxfp4.

The test covers both scaling modes, round_to_input_dtype on and off, bf16 and
fp16 input, plain / transposed / padded / shuffled scales, odd M, and an
all-zero 32-block at each end of the output. In "ceil" mode the Triton scale
clamps amax to 6 * 2^-126, so an all-zero block gets scale 1 where RoundUp
gives 0. The fp4 values of that block are 0 either way, the test expects 1.

Also use device="cuda" in the tests this PR adds, like the rest of the file
since #6783.

Signed-off-by: adlashab <adlashab@amd.com>

---------

Signed-off-by: adlashab <adlashab@amd.com>
Co-authored-by: Shao-Chun Lee <Shao-Chun.Lee@amd.com>

* [CI] Run the selected Triton tests on each merge to main (#7229)

Since #7029, main runs Triton Test only once a night, as the full
suite. A merge that breaks a Triton test then shows up only the next
day, and the nightly cannot say which merge it was.

A push to main now runs Triton Test again, with the same test selection
as a PR: the tests that the merged commit affects, against its parent on
main. MI300X runs only when the selection asks for it, as on a PR. Each
push gets a concurrency group of its own, so a merge no longer cancels
the run of the one before. A push run does not wait for the pre-checks
and does not post on the nightly status issue.

The nightly and manual runs still run the full suite on MI35X and
MI300X.

Co-authored-by: Xin Huang <Xin.Huang@amd.com>

* [FlyDSL] [gfx1250] mega_moe S1F: reachable stage-1 hook, plan precompile, tuned defaults (on by default) (#7315)

* [FlyDSL] [gfx1250] mega_moe S1F: reachable stage-1 hook, plan precompile, opt-in tuning knobs

On top of the stage-1 fusion (#6564):

- Fix: swap the stage-1 hook into dispatch_context as well as scatter. The
  grouped GEMM reads it off dispatch_context, so replacing scatter alone
  left the probe / fused stage-1 path unreachable on the compact path.
- Precompile every stage-1 compact-plan variant at setup (compile only,
  AITER_S1F_PRECOMPILE=0 skips). Eager widths otherwise JIT-compiled inside
  the serving loop: 250-470 ms prefill steps for the first ~100 s of a run.
- AITER_S1F_MAX_RECV=N: steps whose receive bound exceeds N run the whole
  S1F=0 path (plan, GEMM1, no probe), e.g. fuse decode only.
- AITER_S1F_ANY_TILE=1: fuse when the tuned GEMM1's tile_m matches; the
  stage-1 kernel runs its own 256x256 n/k tiles.
- AITER_S1F_PRODUCERS / AITER_S1F_BLOCK_RUNS: per tokens-per-rank bucket
  overrides of the stage-1 tables.
- AITER_TDM_COMPACT_ROUTE_UPTO / _ROUTE_SPEC / _ROUTE_PARALLEL: keep the
  route-parallel compact walk up to N tokens per rank.

All knobs default to the current behaviour.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

* [FlyDSL] [gfx1250] mega_moe S1F: make S1F and its tuned settings the default

Make S1F and its tuned settings the default on gfx1250; env vars become
optional overrides. With nothing set, gfx1250 MegaMoE now runs the
configuration measured on DSV4.1-Flash DEP4 (MI455X, MLPerf Endpoints
agentic c512, +2.1% gen tok/s/GPU over S1F off):

- AITER_MEGA_S1F: on by default; =0 turns the stage-1 fusion off. Configs
  or steps that do not fit still fall back to dispatch + GEMM1; the
  fallback logs at info level unless AITER_MEGA_S1F=1 was set explicitly.
- S1F_BLOCK_RUNS table: (4, 2, 8, 8). AITER_S1F_BLOCK_RUNS /
  AITER_S1F_PRODUCERS still override the tables.
- AITER_S1F_MAX_RECV: 4096 by default; 0 means no limit.
- AITER_S1F_ANY_TILE: on by default; =0 requires all GEMM1 tiles to match.
- AITER_TDM_COMPACT_ROUTE_UPTO / _ROUTE_SPEC: 1024 / 384,16 by default
  (config._COMPACT_ROUTE_UPTO / _SPEC); UPTO=0 keeps _DISPATCH_COMPACT.
- AITER_S1F_PRECOMPILE stays on by default.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

* Read the new S1F/route env vars through the shared env module

Register the env vars this PR adds or changes (AITER_MEGA_S1F,
AITER_S1F_PRECOMPILE / MAX_RECV / ANY_TILE / PRODUCERS / BLOCK_RUNS,
AITER_TDM_COMPACT_ROUTE_UPTO / SPEC / PARALLEL) in mega_moe/envs.py, with
their defaults and parsing, and read them as envs.<NAME> instead of
os.environ. Env vars that were already on main are left as they are.
No behavior change.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

* Drop the new env vars; the tuned S1F/route settings are fixed defaults

This PR no longer adds any environment variable. mega_moe/envs.py is back
to main, and the settings tuned on DSV4.1-Flash DEP4 (MI455X, MLPerf
Endpoints agentic c512) are now plain defaults:

- stage-1 plans are always precompiled when S1F is on;
- _S1F_MAX_RECV = 4096 (s1f_ctl.py);
- GEMM1 fuses whenever its tile_m matches (stage-1 runs its own n/k tiles);
- S1F_PRODUCERS / S1F_BLOCK_RUNS come from the s1f_stage1.py tables
  (block runs 4, 2, 8, 8);
- the compact dispatch always walks routes up to _COMPACT_ROUTE_UPTO = 1024
  tokens/rank with _COMPACT_ROUTE_SPEC = (384, 16) (config.py); the
  AITER_TDM_COMPACT_BLOCKS/WARPS override behaves as on main.

AITER_MEGA_S1F keeps its original os.environ read; only its default
changes to on (0 turns it off). S1F fallbacks log at info level.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

---------

Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>

* [CI] Reuse FlyDSL AOT cache and reduce A8W8 GEMM test overhead (#7316)

* [CI] Reduce A8W8 GEMM benchmark repetition and preserve scale-mode coverage

Signed-off-by: yifehuan <278623243+yifehuan@users.noreply.github.com>

* Use supported MXFP8 row sizes in focused numerical tests

Signed-off-by: yifehuan <yifehuan@local>

* [CI] Reuse FlyDSL AOT cache and sample A8W8 benchmark timing

Signed-off-by: yifehuan <yifehuan@amd.com>

---------

Signed-off-by: yifehuan <278623243+yifehuan@users.noreply.github.com>
Signed-off-by: yifehuan <yifehuan@local>
Signed-off-by: yifehuan <yifehuan@amd.com>
Co-authored-by: yifehuan <278623243+yifehuan@users.noreply.github.com>
Co-authored-by: yifehuan <yifehuan@local>

* [FlyDSL] Fix grouped GEMM C ABI overflow for large EP buffers (#7299)

Pass address-only grouped GEMM operands as typed pointers instead of flattened tensors. FlyDSL packs runtime tensor dimensions as int32, so buffers with more than 2^31-1 elements fail in the generated C ABI fill function. The EP arena handle remains a constexpr and is already excluded from runtime packing.

Remove redundant scale-view normalization and replace its tests with host ABI coverage for large buffers, storage offsets, pitched layouts, cached dispatch, 48-bit addresses, and changed arena specializations.

Signed-off-by: Felix Li <felix.li@amd.com>

* [Triton/Gluon] [HIP] [OPUS] [AITERKER-209] Add gfx1250-strict architecture plumbing (#6767)

* [JIT] Add gfx1250-strict architecture plumbing

* [JIT] Fix ruff F821 and black formatting in gfx1250-strict plumbing

Add the missing build_targets._parse_gpu_archs_env import in core.py
(F821 flagged by CI ruff) and apply black to core.py and
test_build_targets_architecture.py.

Co-Authored-By: Claude Code <noreply@anthropic.com>

* [JIT] Add gfx1250-strict kernel inventory and WS2 build driver

Skeleton inventory for every JIT/FlyDSL/Triton path plus the driver that
fills its strict-build column from a strict-capable container.

Driver verified offline on a non-strict host (module enumeration, child
plumbing, skip classification, --update-md folding):

- CHILD called os.chdir without importing os, so every module would have
  failed with NameError and the whole matrix would have read "fail".
- optCompilerConfig lists 135 modules but get_args_of_build("all") yields
  123. The other 12 are not failures: 10 end in "tune" (skipped by design)
  and 2 are experimental (gated by AITER_ENABLE_EXPERIMENTAL). They are now
  recorded as skip-tune / skip-experimental instead of "missing", so the
  matrix keeps an explicit result for all 135.

Actual pass/fail rows still require the strict toolchain; stock ROCm 7.0.0
rejects gfx1250-strict as an invalid target ID.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

* [JIT] Fix empty GPU_ARCHS raising out of build_module

The gfx1250-strict plumbing routed GPU_ARCHS through
build_targets._parse_gpu_archs_env, which raises when the value holds no
architecture names. GPU_ARCHS set but empty -- what a Dockerfile's
`ENV GPU_ARCHS=` or a bare `export GPU_ARCHS=` produces -- previously just
meant "no cluster-launch flag, build normally"; it now raised RuntimeError
inside build_module and failed every JIT build. Blank, whitespace-only and
";" values had the same problem.

Fall back to the detected arch instead, matching aiter/ops/opus/_arch.py,
which treats an empty or "native" GPU_ARCHS as "probe the device" and
never raises. The arch test moves into _build_archs_from_env /
_cluster_launch_target so it is reachable from a CPU-only test.

Covered by TestClusterLaunchBuildArchs: blank/unset/"native" fall back to
detection, explicit archs win, and -DAITER_ENABLE_CLUSTER_LAUNCH still
follows gfx1250 or gfx1250-strict anywhere in a multi-arch GPU_ARCHS.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

* [JIT] Inventory the gfx1250 dispatch sites a strict device silently misses

The infra PR keeps gfx1250-strict distinct through parsing, build selection
and the asm guard, but Python dispatch still compares against the exact
literal "gfx1250". On A0 get_gfx() returns "gfx1250-strict", so those
comparisons are False and the constant arch allowlists do not contain it:
the gfx1250 path is skipped with no error, which reads as a kernel-selection
change rather than a failure.

ws2_gfx1250_dispatch_sites.py enumerates them -- 150 uses across 65 files,
132 runtime comparisons and 18 allowlist entries -- and records where the
compared string comes from. get_gfx/get_gfx_runtime sites keep the
hyphenated spelling and are known to miss; triton-arch sites read Triton's
own target name, whose strict spelling stays unverified until the toolchain
lands.

The table is generated into the inventory between HTML markers and the
rewrite is idempotent, so it can be refreshed as owners decide each site.
No dispatch behaviour is changed here: every row is an explicit pending
decision, which is what workstream 2 asks for.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

* [JIT] Pin the offload-target parser against hyphenated target names

_so_offload_archs backs the guard that keeps a gfx1250 and a gfx1250-strict
.so from standing in for each other, so the shapes it must parse deserve a
test of their own: a feature list after ':', the strict suffix, both
together, and a ROCm generic target.

The previous gfx\w+ pattern truncated at the first hyphen. That reported a
strict .so as plain gfx1250 -- the substitution this guard exists to
prevent -- and also read gfx9-4-generic as gfx9, which matches no running
arch and so rebuilt such a module on every import.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

* [JIT] Inventory the arch-keyed tables that have no strict entry

The dispatch-site scan only saw comparisons against the "gfx1250" literal,
which left a blind spot: tables keyed by architecture. A strict device is
simply absent from those, so it takes whatever the lookup does with a
missing key.

Four of them carry a gfx1250 entry and no strict entry, and they do not
behave alike:

- radix_topk_one_block._LONG_RADIX_BY_ARCH and ._ONE_BLOCK_LDS_BUDGET_BYTES
  read with a default, so A0 quietly falls back to the generic radix layout
  and a zero LDS budget, losing the 320 KiB one-block path with nothing
  logged.
- pa_decode's inline fp8 dtype table reads with .get() and raises
  NotImplementedError naming the arch -- the "fail clearly" behaviour the
  enablement plan asks for.
- mhc.MHC_FUSED_POST_PRE_M_UPPER_BOUND reads with a default of 1024, which
  happens to equal the gfx1250 entry, so A0 is correct today only by
  coincidence.

The generated table records the lookup style per row so the difference
between a silent fallback and a clear failure is visible to the owner
making the call. Behaviour is unchanged.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

* [JIT] Extend the gfx1250 site inventory to C++ and skip handled sites

get_gpu_arch() in aiter_hip_common.h cuts the target id at ':' and keeps
everything before it, so a strict device reports "gfx1250-strict" in C++
too and the arch_id == "gfx1250" branches miss exactly as the Python ones
do. 27 of them across 14 files, 11 in mha_bwd.cu alone, where several pick
launch geometry rather than a code path -- a miss changes a block size
instead of raising.

Worth recording for the Opus owner: opus_co_launch_gfx1250.cuh builds its
code-object path from get_gpu_arch(), so A0 looks for gen_co/gfx1250-strict
and fails with a clear AITER_CHECK rather than loading a B0 object. That is
what the plan asks for, so it is called out in the table's preamble instead
of being listed as a gap.

The scan now skips any site with a strict spelling within four lines, so
the asm gate's own early return no longer pads the table and a site drops
out once an owner has handled it.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

* [JIT] Count the tuned rows a strict device cannot use

Tuned GEMM and MoE configs carry a gfx column and are looked up through
get_gfx(), which reports gfx1250-strict on A0. 1620 rows across 18 config
files are keyed to gfx1250 and none to gfx1250-strict, so on A0 every shape
they cover falls back to the default path.

That is what the plan asks for -- strict tuning has to be explicit and
measured, not a filename alias -- so the generated table is the workstream
3 worklist rather than a defect list: each file needs either measured
strict rows or a recorded decision that the B0 rows hold on A0.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

* [JIT] Generate the module table's gfx1250 signal column

The column was hand-filled and followed no single rule: 13 of 135 rows
disagreed with any consistent reading of the sources, in both directions.
module_custom_all_reduce_gfx1250 was marked by its module name while
module_cache was marked `guard` although its sources also compare against
the "gfx1250" string literal. Owners use this column to decide where to
look, so wrong values send them to the wrong file.

--refresh-signals now computes it from the sources that
core.get_args_of_build actually resolves, falling back to the raw config
expressions for the tune and experimental modules that call skips. The
markers are explicit and can combine: `name` for the module name, `str` for
the "gfx1250" string literal, `macro` for the __gfx1250__ guard.

The distinction matters for A0 risk. `macro` rows keep compiling, because
the strict target defines both __gfx1250_strict__ and __gfx1250__. `str`
rows are the runtime decisions a strict device silently misses. The refresh
changed 27 cells and is idempotent.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

* [JIT] Tie the three arch tables together with a test

Adding a target means editing GFX_MAP, GFX_CU_NUM_MAP and core.py's
allowed_archs, and nothing tied them together. allowed_archs carried a
duplicate "gfx1250" entry until this series removed it, which is what that
drift looks like in practice.

Two invariants hold today and are now pinned: every arch with an offline CU
default is accepted by validate_and_update_archs and maps back through
GFX_MAP, and gfx1250-strict does the same. Deleting gfx1250-strict from
allowed_archs fails both tests, so a half-landed target is caught in CPU-only
CI rather than at build time.

The check is behavioural -- it drives validate_and_update_archs and
get_gfx_custom_op_core with GPU_ARCHS set -- so it also covers the arch
reaching hipcc unchanged as --offload-arch, hyphen included.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

* [JIT] Record CK modules as skipped when the submodule is not checked out

aiter's CK-free auto-exclusion tests os.path.isdir(CK_3RDPARTY_DIR), which
is True for an uninitialised submodule: the directory exists, it is simply
empty. A checkout without `git submodule update` therefore still tries to
build the CK modules, and they fail on missing sources.

core._get_ck_exclude_modules() reports 69 of the 135 modules as
CK-dependent, so on such a checkout more than half the build matrix would
fill with failures that say nothing about the strict target. The driver now
treats an empty composable_kernel directory as CK-free and records those
modules as skip-no-ck, keeping pass/fail for results that are actually
about gfx1250-strict.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

* [JIT] Pin PYTORCH_ROCM_ARCH in the build driver

Setting GPU_ARCHS alone does not reach the compiler. torch's cpp_extension
takes --offload-arch from PYTORCH_ROCM_ARCH, so with GPU_ARCHS=gfx1250 the
modules still compiled for torch's own default list -- observed in a ROCm
10.0.0 container as

  --offload-arch=gfx90a ... --offload-arch=gfx1151

with no gfx1250 anywhere. Every row of the matrix would have been a result
for the wrong architectures while looking like a validated one. The driver
now derives PYTORCH_ROCM_ARCH from GPU_ARCHS; with that fixed,
module_aiter_unary builds for gfx1250 and build.ninja carries exactly one
--offload-arch.

Failures also kept the wrong 15 lines: ninja echoes the whole compile
command, which pushes the diagnostic out of a plain tail, so a failing
module recorded only the Python traceback. The recorded tail now leads with
up to five "error:" lines.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

* [JIT] Record what ROCm 10.0.0 does and does not provide for strict

Checked in rocm/atom-dev:latest-rocm10.0 -- ROCm 10.0.0, AMD clang
23.0.0git (ROCm/llvm-project 8f497e0992fb), HIP 7.15.26333. hipcc compiles
gfx942 and gfx1250 and rejects gfx1250-strict as an unsupported HIP gpu
architecture. The processor list is gfx1250 and gfx1251, and neither
libLLVM nor amdclang contains the string gfx1250-strict or
__gfx1250_strict__ at all.

The only A0/B0 notion in that toolchain is the AMDGPU subtarget feature
gfx1250-b0-specific, and it is not reachable from the driver: not a
processor, rejected inside a target id, and ignored as a clang target
feature.

So the build matrix cannot be filled from the public ROCm 10 container --
the plan's toolchain dependency is unmet rather than satisfied -- and there
is a question worth settling before the identity model is locked in. If the
mechanism that ships is a subtarget feature on gfx1250 rather than a
separate processor name, an A0 and a B0 object both carry the target id
gfx1250, which is exactly what _so_offload_archs reads, so that guard would
need a different discriminator. The survey records the evidence and the
question without presuming the answer.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

* [JIT] Lead the recorded failure tail with the real compiler diagnostic

The error capture grabbed the first five lines containing "error:", which
on a failing build are usually the Python RuntimeError and
CalledProcessError wrapper lines, not the compiler's own message. Prefer
"file:line:col: error:" frontend diagnostics when present, fall back to any
"error:" line otherwise (backend errors such as the gfx1250 DPP rejection
have no file:line prefix), and de-duplicate -- clang repeats each
diagnostic once per offload arch.

Verified against a real ROCm 10 matrix run: the ck_tile "extraneous closing
brace" / "template exec" and ck "operator[] in Tuple<>" errors now lead
their tails instead of the wrapper, and the opus DPP "broadcasts are not
supported on GFX10+" backend error still surfaces via the fallback.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* [JIT] Record the measured gfx1250 build matrix on ROCm 10.0.0

Ran the full 135-module matrix in rocm/atom-dev:latest-rocm10.0 with
GPU_ARCHS/PYTORCH_ROCM_ARCH=gfx1250: 102 pass, 21 fail, 12 skip. Strict
itself does not compile in this toolchain, so this is the gfx1250 baseline
-- anything that fails for gfx1250 cannot build for strict, which is gfx1250
plus an A0 code-object identity.

Every failure was rebuilt for gfx942 as a control, which corrected my first
read of them. They are not a ROCm-10/clang-23 regression: 18 of 21 build
cleanly on gfx942 and fail only on gfx1250, and the 3 FMHA-bwd modules that
timed out on gfx942 were compiling without error (module_mha_bwd from the
same family finished and passed). Every failure is a genuine gfx1250
porting gap:

- DPP broadcast dpp_ctrl modes GFX12 removed (moe_sorting_opus,
  moe_mxfp4_aux)
- MFMA / buffer-load-LDS intrinsics GFX12 lacks (iq2r_moe,
  mha_fwd_native_splitkv)
- a wavefrontsize64 request on a wave32 part (chunk_gdr_fwd_h)
- warp-size and gfx950-trait static assertions (topk_plain, moe_opus)
- CK's ck_tile TDM path, selected on gfx1250 by
  CK_TILE_USE_BUFFER_ADDRESSING_BUILTIN=0 (the cktile/blockscale GEMMs,
  activation, sample, the ck_tile FMHA fwd/prefill)
- fmha-v3-bwd asm configs not generated for gfx1250 (the mha_bwd family)

The inventory now carries the result table with the gfx942 control column
and the per-module cause, as the workstream-2 porting worklist.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* [JIT] Correct the strict macro model and inventory __gfx1250__ guards

The plan claimed the strict target defines both __gfx1250_strict__ and
__gfx1250__, so existing __gfx1250__ guards keep compiling. ROCm 10.1's
actual behaviour is the opposite: strict defines __gfx1250_strict__ and
does NOT define __gfx1250__ (ROCm/rocm-systems#12068). A __gfx1250__-only
device guard therefore compiles its body away on a strict agent -- a silent
empty kernel, the upstream TDM_Basic_load_2d failure, with no build error.
JAX hit the matching compiler-side fact and dropped strict from its ROCm
10.0 defaults (jax-ml/jax#41187), and ROCm 10.1 adds gfx1250-strict as a
distinct --offload-arch target across the stack (RCCL, CK, rocSHMEM), which
confirms this PR's distinct-name identity model and the .so guard that keys
on it.

Fixes, docs only -- no source guard is touched, because the change must not
break the shipping ROCm 7 build, whose compiler knows only __gfx1250__ and
has no strict target:

- plan and inventory legend now state the real macro behaviour and the
  backward-compatible fix, `defined(__gfx1250__) || defined(__gfx1250_strict__)`,
  whose strict arm is simply never taken on ROCm 7.
- the ROCm 10.0 survey's open "name vs feature" question is resolved: 10.1
  ships a distinct target name; the 10.0 gfx1250-b0-specific subtarget
  feature was an early-LLVM artifact.
- ws2_gfx1250_dispatch_sites.py now also inventories the __gfx1250__
  preprocessor guards (95 across 24 csrc files) as a generated review table,
  flagging which already combine with ||/&& to read first. This is the class
  the earlier string-comparison scan did not cover.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* [JIT] Make every __gfx1250__ device guard also fire on strict

Non-asm operators are meant to run on both MI450 SKUs (A0 = gfx1250-strict,
B0 = gfx1250); only the shipped asm is B0-only. But the strict compiler
defines __gfx1250_strict__ and not __gfx1250__, so every
`#if defined(__gfx1250__)` device guard took its #else branch on a strict
build -- a silent wrong path that still compiles. Read on A0 this is:
get_smem_size returning 64 KB instead of 320 KB, mhc falling back to fp32
instead of the WMMA path, make_gmem returning the 0xffffffff sentinel so
stores drop, and the TDM sync counters using the legacy lgkmcnt/vmcnt
instead of the split dscnt/loadcnt.

Rewrite all 97 `defined(__gfx1250__)` occurrences across 24 csrc files to
`(defined(__gfx1250__) || defined(__gfx1250_strict__))`. The parentheses
keep it correct in every context: positive guards, `!(...)` exclusion
lists (strict is excluded exactly as gfx1250 is), and `||`/`&&` family
lists. The strict arm is never taken on the shipping ROCm 7 build, which
has no strict target, so that build is unaffected.

Compile-verified on ROCm 10.1 (clang 24.0.0git) for both gfx1250-strict and
gfx1250 on module_mhc and module_top_k_per_row; both pass for both targets
at low parallelism (the earlier full-matrix failures for these were -j179
contention, not real). Runtime correctness on A0 is left to the daily A0
op-test once the modules are built as gfx1250-strict.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* [JIT] Add strict to the compiled-path gfx1250 dispatch comparisons

Non-asm operators run on both MI450 SKUs, and strict is gfx1250 (A0)
hardware, so a `get_gfx() == "gfx1250"` dispatch that selects a compiled
path must also fire on gfx1250-strict -- otherwise A0 silently takes the
generic fallback and, where the gfx1250 branch picks an arch-specific
layout (wave32 tiling, LDS budget), gets a wrong result, not just a slow
one.

Rewrite the 101 compiled-path comparisons across 51 files from
`== "gfx1250"` / `in ("gfx1250",)` to `in ("gfx1250", "gfx1250-strict")`.

Deliberately NOT changed: the one asm-selecting dispatch,
gemm_op_a4w4.py `if gfx_arch in ["gfx1250"]`. Adding strict there would
route A0 into the f4gemm asm path, which the asm gate then throws on --
and the policy is that asm is simply unavailable on A0, so strict must keep
missing that comparison and fall to the compiled path. The classifier that
split compiled from asm sites flagged exactly this one site.

py_compile, ruff and black clean (topk.py's pre-existing fmt:skip parse
issue aside); the CPU-only arch/dispatch tests still pass. Multi-element
arch allowlists (gluon_archs, SUPPORTED_ARCHS) are left for a follow-up --
those are a capability/feature gap on strict, not the correctness-level
layout divergence this commit closes.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* [JIT] Cover != "gfx1250" and post-merge dispatch sites for strict

The first dispatch pass only rewrote == and single-element `in (...)`
comparisons, so `!= "gfx1250"` negations were left behind -- on strict
`arch != "gfx1250"` is True and the kernel takes the non-gfx1250 branch,
the same silent-wrong-path class the == sites had. The main merge also
added fresh gfx1250 comparisons (flydsl pa_decode, triton quant/mla,
mla backend plan).

Fix the 30 compiled-path sites: `!= "gfx1250"` ->
`not in ("gfx1250", "gfx1250-strict")`, `== "gfx1250"` -> `in (...)`.
Three asm-selecting sites are deliberately left (attention.py:1752,
mha.py:3779, mla_sparse_prefill.py:100) so strict keeps missing them and
A0 errors on asm rather than loading it, per the policy that asm is
unavailable on A0.

py_compile, ruff and black clean; CPU-only arch tests pass. Only the three
asm sites now compare against the bare "gfx1250" literal.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* [JIT] Add strict to the C++ arch_id == "gfx1250" dispatch comparisons

get_gpu_arch() in aiter_hip_common.h keeps the full target id, so a strict
device reports "gfx1250-strict" and every `arch_id == "gfx1250"` branch in
csrc misses -- the same silent-wrong-path class the Python and __gfx1250__
guards had. has_tdm_arch() returned false on strict even though A0 has TDM;
mha_bwd picked the wrong launch geometry; and so on.

Rewrite the 28 comparisons across 14 files: `X == "gfx1250"` ->
`(X == "gfx1250" || X == "gfx1250-strict")`, `X != "gfx1250"` ->
`(X != "gfx1250" && X != "gfx1250-strict")`, where X is a plain identifier
or get_gpu_arch(). The asm-dispatch sites are included too: the actual asm
load still funnels through require_gfx1250_asm_or_throw, which throws on A0,
so computing the dispatch as gfx1250 is safe and A0 still cannot load asm.

The already-strict-aware asm gate in aiter_hip_common.h is left untouched.
Paren balance verified; the gfx942/gfx950 C++ build in CI compiles it.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* [JIT] Add strict to the multi-element arch allowlists

The capability allowlists -- gluon_archs, SUPPORTED_ARCHS, _SUPPORTED_GFX,
DEVICE_ARCH tuples, the triton tuning kernel decorators -- list gfx1250 as
a supported arch but not gfx1250-strict, so strict (which is gfx1250
hardware) was treated as unsupported: no gluon path, excluded from the
tuned kernel sets, etc.

Add "gfx1250-strict" next to each "gfx1250" element in those arch tuples
and lists (41 across 22 files), skipping the substring-match helper
(quick_all_reduce's `gfx in arch`, where "gfx1250" already matches
"gfx1250-strict") and any asm-selecting context. With the gemm_a4w4 path
main restructured, the gfx1250 branch now routes plain MXFP4 to FlyDSL and
the rest to asm, so strict correctly gets the FlyDSL path and the asm gate
still handles A0.

ruff and black clean; CPU-only arch tests pass. This closes the last
aiter-side strict dispatch gap; what remains is A0 runtime validation and
the Opus .co rebuild.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* [JIT] Pin flydsl to 0.4.0.dev973 for gfx1250-strict LDS support

FlyDSL's arch handling is otherwise strict-ready because it matches by
prefix (arch.startswith("gfx1250") / "gfx12"), which gfx1250-strict
satisfies -- WMMA instruction selection, warp_size (32) and the compile
target all come out right. The one gap was the LDS-capacity map, keyed on
the exact "gfx1250" string, so on strict the lookup missed and
check_smem_capacity silently skipped the overflow guard.

0.4.0.dev973 adds "gfx1250-strict": 327680 to SMEM_CAPACITY_MAP (now in
flydsl.runtime.device), closing that gap. Bump the pin in requirements.txt
and setup.py (the version assertion and auto-install both derive from
FLYDSL_VERSION). Revert to the default release once the fix ships there.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* [JIT] Clean up: drop investigation tooling, tidy strict macro guards

Self-review cleanup, no behaviour change:

- Remove the ws2_* enumeration/matrix scripts and the generated
  kernel-inventory doc. Those were one-off investigation flow (build
  matrix, dispatch-site enumeration) and do not belong in the repo; nothing
  in aiter imports them.
- Drop the dangling "former hotswap path" bullet from the enablement plan;
  the term exists nowhere in the tree.
- Flatten the redundant parentheses the strict sweep left behind: the
  mechanical `defined(__gfx1250__)` -> `(defined(__gfx1250__) ||
  defined(__gfx1250_strict__))` wrapped the pair even inside OR chains.
  Unwrap it where the context is pure-OR (90 sites); keep the parentheses
  only where they are load-bearing -- negation `!(...)` and `&&` operands
  (7 sites). Compile-verified on ROCm 10.1 for both gfx1250-strict and
  gfx1250 (module_mhc, module_top_k_per_row); CPU-only arch tests pass.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* fix: pin flydsl build-system dep to 0.4.0.dev973 to match setup.py

pyproject.toml build-system.requires still pinned flydsl==0.4.0 while
setup.py (FLYDSL_VERSION) and requirements.txt use 0.4.0.dev973. Under
build isolation pip installed 0.4.0 as the build dep, then setup.py's
version assertion rejected it and the build failed. Align all three.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* fix: asm-gate error message covers both A0 detection routes

The gate blocks B0 asm on A0 via two routes: gcnArchName == gfx1250-strict,
or gcnArchName == gfx1250 with asicRevision 0. The message hardcoded
'gfx1250-strict/A0', which misleads when A0 is caught by the asicRevision
route (device name is literally gfx1250, rev 0). State both forms.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* fix: strict Opus codegen + 4 CI-red CPU tests

gen_instances: GPU_ARCHS=gfx1250-strict looked up KERNEL_REGISTRY by the
strict name, which has no registry key, so Opus codegen emitted an empty
header and module_deepgemm_opus failed to build (the only strict-specific
failure of the strict build matrix). Fold the registry key to the gfx1250
family (mirroring chip_info._detect_native) while keeping the gfx1250-strict
build target, so hipcc still compiles for mach 0xeb. Strict now resolves to
256 CU, same as gfx1250.

Tests (Draft CI had skipped these):
- test_opus_co_profile load_chip_parsers: keep the _GFX_ARCH_PATTERN constant
  in the extracted AST namespace (the strict PR added it; parsers NameError'd).
- test_build_targets preserve->normalize: _detect_native now folds strict to
  gfx1250 for dispatch; assert the new behavior.
- test_launcher_does_not_pin_the_first_device: use the first declared local
  kid instead of hard-coded global kid 21016 (removed by #6937).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* fix: strict Opus codegen emits the gfx1250 family (routes>0), not an empty module

The earlier fix only folded the target list; the real route generation still
queried the registry by the strict name and produced routes=0 (empty
module_deepgemm_opus). Apply the gfx1250 family name at every registry/namespace
lookup -- available's 'if gfx in arches', get_registry, default_compiled_keys --
and at the OPUS_BUILD_HAS_* macro (C checks #ifdef OPUS_BUILD_HAS_GFX1250, and a
hyphenated STRICT macro is invalid C). KernelKey and the build target keep the
gfx1250-strict spelling, so hipcc still offloads for mach 0xeb.

Verified: GPU_ARCHS=gfx1250-strict now yields routes=228 (== gfx1250) with
OPUS_BUILD_HAS_GFX1250; CPU tests stay green.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* [JIT] Opus codegen: normalize gfx1250-strict targets to the gfx1250 family at the source

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* fix: bump flydsl pin to 0.4.0 to match main (dev973 would downgrade it on merge)

main pins flydsl==0.4.0; this PR still carried the 0.4.0.dev973 prerelease,
so merging would silently downgrade the whole repo. 0.4.0 is a superset of
dev973 (contains the gfx1250-strict LDS fix) and is what main already uses.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* test: drop redundant test_strict_resolves_through_all_three_tables

test_every_offline_target_resolves_through_all_three_tables already loops
over GFX_CU_NUM_MAP, which contains gfx1250-strict, so the strict-specific
case was an exact duplicate of one loop iteration.

* fix: avoid 'in' tuple test in gluon kernel constexpr (unsupported by gluon)

The strict-allowlist sweep turned 'ARCH == "gfx1250"' into
'ARCH in ("gfx1250", "gfx1250-strict")' inside gluon kernel bodies/constexpr.
Triton's gluon compiler does not implement the 'In' AST operator, so these
kernels failed to compile ('AST comparison operator In is not implemented'),
breaking test_pa_mqa_logits_2d_context_lens[gluon*] on gfx942/gfx950. Use
'(x == a or x == b)', which gluon supports, in kernel scope. Host-side Python
'in' checks are unaffected and left as-is.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* Preserve exact gfx1250-strict identity; map to gfx1250 family only at dispatch

Per the ROCr rollout contract (all SW consumes the exact GFX ID ROCr reports),
global normalization in _detect_native() is the wrong boundary. Once ROCr
reports gfx1250-strict for A0, folding it to gfx1250 there makes
_needs_arch_rebuild() treat a strict .so as not covering a strict device and
rebuild in a loop; it also remixes A0/B0 in tuned configs, caches, and logs.

- chip_info._detect_native: drop the strict->gfx1250 fold; return the exact
  ROCr-reported identity. get_gfx_runtime()/get_gfx() now report gfx1250-strict
  on A0.
- build_targets.py: add the single source of truth for family folding --
  GFX_BASE_ARCH, base_gfx_arch(), is_gfx1250_family(). Torch-free, so build
  scripts and tests import it without side effects.
- Opus: replace its two private strict->gfx1250 folds (gen_instances
  _registry_arch, opus_gemm_common _get_registry_records/get_registry/
  default_compiled_keys) with base_gfx_arch(), so there is no second mapping.
  hipcc still gets --offload-arch=gfx1250-strict and the .co stays mach-0xeb.
- mega_moe _require_gfx1250: restore the gfx1250-family guard (via
  is_gfx1250_family) that main's #6430 refactor collapsed to bare gfx1250.
- Build/cache/tuning/.so/.co/ASM keep the exact identity: _needs_arch_rebuild
  and _so_offload_archs already key on it, so A0 never loads a B0 object.
- Pre-existing gfx1250-only dispatch sites authored by other PRs stay exact;
  with the fold gone they correctly distinguish strict (fall back) instead of
  silently reusing a B0 path.
- Tests: flip the native-detection test into a regression lock that strict
  identity is preserved; add family-helper and 4-way code-object rebuild-guard
  tests (runtime/built strict x gfx1250 -> accept/reject).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* Apply ruff import sort and black formatting to the identity refactor

Fixes the pre-checks (ruff I001 import order in mega_moe, black reformat of
mega_moe and test_build_targets_architecture) that gated check-signal and
skipped the Triton/Aiter test runs. No logic change.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* black: wrap the long gfx1250/gfx1250-strict dispatch conditions in the gluon kernels

The in->or fix in the pa_mqa_logits / unified_attention_2d gluon kernels left
two over-length boolean conditions that Black wraps. Formatting only; the
gluon compiler still sees the same (x == a or x == b) expression.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* Define Opus base_gfx_arch locally instead of importing build_targets

The cross-package `from build_targets import base_gfx_arch` (via a
__file__-relative sys.path insert) breaks the Aiter wheel build:
gen_instances.py runs from a staging tree (aiter_meta/csrc/opus_gemm) where
../../aiter/jit/utils does not exist, so module_deepgemm_opus failed with
ModuleNotFoundError: No module named 'build_targets'. Duplicate the one-line
helper locally in opus_gemm_common (documented to mirror
build_targets.base_gfx_arch) so the standalone Opus codegen has no cross-tree
import dependency. gen_instances keeps importing it from opus_gemm_common.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* test(jit_cache): seed _cluster_launch_target in the build_module sandbox

test_jit_cache_transaction AST-extracts build_module into a sandbox namespace
with CPU deps injected. Since the gfx1250-strict work added a
_cluster_launch_target(get_gfx) call inside MainFunc, the sandbox hit
NameError: name '_cluster_launch_target' is not defined (25 TestBuildPublication
errors on MI35X). This surfaced only now because the earlier Opus wheel-build
failure skipped the standard tests. Inject the dependency like the others; for
the gfx942 test context the real function returns False, so the mock matches.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

---------

Co-authored-by: zufayu <8138883+zufayu@users.noreply.github.com>
Co-authored-by: Claude Code <noreply@anthropic.com>
Co-authored-by: root <root@hjbog-srdc-24.amd.com>

* [HIP] [FlyDSL] optimize dsv41 perf for decode shape (#7309)

* optimize the mxfp8 op for decode shape

* retune mxfp8 tables with optimize decode shape

---------

Co-authored-by: solin <bingzhou@amd.com>

* [HIP] [CK] [JIT] [Tuner] Rank accuracy-qualified kernels by package energy (#6970)

* [Tuner] Rank accuracy-qualified kernels by package energy

Add --select power so shared tuners keep the lowest runtime times PPT watts. The accuracy gate is unchanged, and latency remains the default.

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

* [Tuner] Satisfy black and ruff on the power-selection change

Format the tuner edits and replace the NaN self-comparisons and bare exception handlers that the pre-check jobs reject.

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

* [Tuner] Keep latency order when package power is missing

Sort power candidates by energy and then latency so an unread sensor cannot drop the fastest kernel. A single AITER_POWER_PATH is rejected when more than one tuner process is running.

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

* [Tuner] Keep power-tuned rows apart from latency tables

Power tunes of shipped tables now write to aiter/configs/power_configs/, which loads only with AITER_SELECT=power and replaces only the shapes it covers. Each power row records perf level, power cap, and SCLK range, and the config merge refuses to auto-resolve power rows against latency rows or rows tuned under different conditions. Power mode is marked experimental.

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

* [Tuner] Isolate power-kernel selection from the tuner and JIT path

Ranking, package-power measurement, and power_configs loading each have one owner. Latency CSV columns stay unchanged unless --select power.

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

* [Tuner] Measure power with the same call the timer uses

The replay was passing run_perftest keywords into the kernel, so watts came back NaN and the row was still stamped select=power. A batch with no finite watts stays a latency result, and --select power is rejected on tuners that do not record it.

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

* [Tuner] Give the MXFP4 power retime test a supported shape

Main now rejects an MXFP4 row before the scan unless dtype and quant fields match the port. The energy-ranking fixture was missing those fields, so it never reached the retime.

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

* [Tuner] Format the power-selection files with Black

CI skipped the GPU tests because Black rejected these three files.

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

* [Tuner] Keep the wheel build from importing the aiter package

setup.py loads jit.core before aiter exists. Importing the power overlay from that path ran aiter/__init__.py while module_aiter_core was still being registered.

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

* [Tuner] Decide latency or power in one policy and measure in one place

A tuner builds a SelectionPolicy from --select and asks it for the
candidate rank, frame order, pair energy, accurate-pass count, and the
power sample. Tuners no longer branch on the select mode or write
AITER_SELECT; that variable only picks the tables jit/core loads.
mp_tuner reads the policy TunerCommon.run activates.

package_power.measure_kernel is the one power read. mp_tuner hands it
the kernel launch with the run_perftest options removed, so the window
replays the kernel back to back. Replaying run_perftest itself left the
GPU idle for most of the window.

The duplicate-shape dedup in jit/core sorts by lowest us without
importing the aiter package when setup.py loads it as jit.core.

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

* [Tuner] Fix power-tune paths found by end-to-end runs on MI355X

- TunerCommon.run wrapped its body in _run, which IQ2R already defines
  as its kernel launcher, so every IQ2R tune failed. Use _run_selected.
- The a16w16 tuner built gradlib's Gemm without the policy, so hipBLASLt
  was never power-measured and always ranked last. Pass the policy, and
  default Gemm and the K5 sweep to the active policy.
- The power window queued launches without bound. IQ2R graph replays
  crash inside the HIP runtime when hundreds are queued. Keep a few
  events in flight, one per ~200 us of work, so the GPU never drains.
- IQ2R times a captured CUDA graph; sample power from that graph too.
- Staged MXFP4 logs each pair's watts like the other loops.

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

* [Tuner] Correct the power sample for host launch gaps

A kernel whose Python launch takes longer than its GPU time leaves the
GPU idle between launches, so the window read close to idle and the
kernel looked efficient (split-K FlyDSL GEMMs at 256x4096x4096 read
~370 W against ~990 W replayed from a CUDA graph, and won power mode).

Every measurement site now passes the call's timed us. When the loop
never waited on the GPU, the window's busy fraction is launches * us /
window, and the reading is scaled to active power over the between-launch
baseline (a one-element kernel loop, ~300 W on MI355X). A window that
waited on the GPU kept it saturated and stands as read. Under 20% busy
the candidate stays unmeasured.

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

* [Tuner] Keep latency profiles as main and write only measured power rows

- MoE latency tunes write the -o2 profile and profile_fmoe.csv with
  main's columns and row order. Stage watts, rank, energy, and
  tflops_per_w appear only with --select power.
- A power tune drops rows whose watts were not measured instead of
  writ…
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants