Skip to content

[CK] [OPUS] [CI] Mi300a enablement - #5186

Merged
afanfa merged 21 commits into
ROCm:mainfrom
afanfa:mi300a-enablement
Nov 11, 2026
Merged

afanfa merged 21 commits into
ROCm:mainfrom
afanfa:mi300a-enablement

Conversation

@afanfa

@afanfa afanfa commented Sep 1, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

Add initial support for MI300A

Technical Details

This PR has three parts: 1) improve the device detection for MI300A; 2) add encoding for fp8 e4m3fnuz ; 3) Add MI300A (gfx942, cu_num=228) configs for GEMM and fused MoE

Test Plan

New tests for the newly added features have been added. The test suite has been run

Test Result

No regressions or failures on the newly added tests.

PR Self-Review

Reviewed revision

518476e

Behavior and intent

Adds MI300A (gfx942, 228 CU) support. Before this PR, MI300A was reported as "MI300" (same as MI300X), so SKU-specific dispatch could not tell them apart. Its 228-CU tuned rows did not exist, so every GEMM/MoE shape fell back to heuristic kernels. Tuning shape lists that name torch.float8_e4m3fn also aborted with "Unsupported dtype" (or were silently skipped by the bpreshuffle tuner) on gfx942, which only implements e4m3fnuz.

  • chip_info: map PCI chip id 0x74A0 to "MI300A" in get_device_name().
  • dtypes.normalize_fp8_dtype() + TunerCommon._normalize_fp8_columns(): resolve either e4m3 flavor in untuned shape lists to the arch-native encoding on ingest (name lookup, no eval), in the shared loader so every tuner family gets it.
  • Tuned configs: MI300A rows for bf16, a8w8, a8w8 blockscale GEMM and fused MoE.
  • Build targets: after merging main, this PR uses upstream's GFX_CU_NUMS_MAP (gfx942 → 304/228/80) unchanged. The earlier live-GPU CU refinement was removed, so wheels do not depend on the build host.

Validation and unverified scope

Hardware: AMD Instinct MI300A (gfx942, 228 CU), ROCm 7.14, torch 2.14.0+rocm7.14, Python 3.11.

  • python -m pytest op_tests/tuning_tests/test_arch_resolution.py op_tests/test_build_targets_architecture.py tests/test_chip_info_xcds.py → 84 passed (this exercises the fn→fnuz direction, since the host is gfx942).
  • python -m pytest op_tests/test_opus_a16w16_policy_parity.py → 19 passed.
  • Level 0+1 suite from tuning-tests.yaml → all pass except 4 tests (test_tuner_infra GroupedFmoe x2, test_fast_tuning a16w16 policy, test_mp_tuner_logic timing) that fail identically on a clean upstream/main checkout in this environment (Triton version), so they are unrelated.

Unverified: the new tuned rows were not re-benchmarked at this revision; the gfx950 (fnuz→fn) direction is covered only by the CI tuning job on MI35X. The Aiter Test wheel build/GPU tests have not completed (Docker Hub 429 on the fork PR's anonymous image pull).

PR-specific question

The only CI change adds op_tests.tuning_tests.test_arch_resolution to the existing Level 0+1 unittest command in .github/workflows/tuning-tests.yaml. It adds no new triggers, permissions, secrets or steps. Fork PR code already runs in that same unprivileged container job, and no PR text is interpolated into a command.

Changed op tests

  • op_tests/tuning_tests/test_arch_resolution.py (new): pins MI300A SKU naming by PCI id, the version-gated PciChipId ordinal, legacy cu_num→gfx backfill, that GPU_ARCHS builds include (gfx942, 228) alongside 304/80, fp8 e4m3 resolution, that every shipped shape list (incl. bpreshuffle) is tunable on the running arch, and that the cu_num=228 tuned rows resolve to gfx942. No existing test covers MI300A identity or fp8 flavor resolution of shape lists.
  • op_tests/test_opus_a16w16_policy_parity.py: the pre-PR codegen reference map now keys on cu_num. Without it, the MI300A (228) and MI308X (80) rows collide with MI300X (304) rows for the same shape and the parity assertion compares the wrong kernel ids. The test is otherwise unchanged.
  • op_tests/tuning_tests/README.md: documentation only; lists the new module and adds it to the suggested Level 1 command.

@afanfa
afanfa requested review from a team and a lite review from Copilot September 1, 2026 15:56
@github-actions

github-actions Bot commented Sep 1, 2026

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; main branch always runs both MI35X and MI300X
multigpu 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: GPT-OSS-120B, DeepSeek-R1-0528, Kimi-K2.5
ci:all All standard extended tests (excludes ci:atom_full)

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

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 opt this PR out.

@github-actions github-actions Bot changed the title Mi300a enablement [CK] [CI] [JIT] Mi300a enablement Sep 1, 2026

Copilot AI 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.

Pull request overview

Adds initial MI300A (gfx942, 228 CU) enablement by improving SKU/CU resolution, normalizing FP8 e4m3 flavor handling (fn vs fnuz), and introducing MI300A-tuned GEMM/MoE configuration rows, with accompanying unit tests wired into CI.

Changes:

  • Refine device/SKU detection and build-target CU selection for gfx942 variants (MI300X/MI300A/MI308X), including a safer PCI chip-id probe path.
  • Add FP8 e4m3 flavor normalization and apply it during tuning/shape-list ingestion so shipped shape lists remain tunable on gfx942/gfx950.
  • Add MI300A tuned configs (GEMM + fused MoE) and add Level-1 unit tests + CI hooks validating arch/dtype resolution invariants.

Reviewed changes

Copilot reviewed 13 out of 13 changed files in this pull request and generated 4 comments.

Show a summary per file
File Description
op_tests/tuning_tests/test_arch_resolution.py New unit tests for gfx942 SKU resolution, PCI chip-id probing behavior, CU-based arch backfill, build-target CU refinement, and FP8 e4m3 normalization.
op_tests/tuning_tests/README.md Documents the new Level-1 test module and adds it to the suggested unittest invocation list.
docs/autotuning_pipeline.md Adds the new test module to the documented Level-1 pipeline list.
csrc/ck_gemm_moe_2stages_codegen/gemm_moe_tune.py Normalizes fp8 e4m3 dtype strings on untuned-shape ingestion; normalizes q dtypes when reading/tuning configs.
csrc/ck_gemm_a8w8/gemm_a8w8_tune.py Normalizes fp8 e4m3 dtype strings for q_dtype_w on untuned-shape ingestion.
aiter/utility/dtypes.py Introduces normalize_fp8_dtype() to map e4m3 flavors to the arch-native fp8 encoding.
aiter/jit/utils/chip_info.py Adds MI300A chip-id mapping, updates legacy CU→gfx backfill, refines get_build_targets() CU selection under GPU_ARCHS, and makes PCI chip-id probing degrade to None.
aiter/jit/utils/build_targets.py Comment updates for gfx942 default CU count context.
aiter/configs/model_configs/mi300a_tuned_fmoe.csv Adds MI300A (gfx942, 228 CU) tuned fused-MoE rows.
aiter/configs/bf16_tuned_gemm.csv Adds MI300A (gfx942, 228 CU) tuned bf16 GEMM rows.
aiter/configs/a8w8_tuned_gemm.csv Adds MI300A (gfx942, 228 CU) tuned a8w8 GEMM rows (incl. fp8 fnuz).
aiter/configs/a8w8_blockscale_tuned_gemm.csv Adds MI300A (gfx942, 228 CU) tuned a8w8 blockscale GEMM rows.
.github/workflows/tuning-tests.yaml Runs the new Level-1 arch-resolution test in CI.

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

Comment thread csrc/ck_gemm_a8w8/gemm_a8w8_tune.py Resolved
Comment thread aiter/jit/utils/chip_info.py Resolved
Comment thread aiter/jit/utils/build_targets.py Resolved
Copilot AI review requested due to automatic review settings September 1, 2026 17:01

Copilot AI 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.

Pull request overview

Copilot reviewed 13 out of 13 changed files in this pull request and generated 1 comment.

Suppressed comments (1)

aiter/jit/utils/chip_info.py:463

  • _get_pci_chip_id() only catches OSError/ValueError/UnicodeDecodeError. If libamdhip64.so loads but does not export hipDeviceGetPCIBusId (or the symbol name changes), ctypes will raise AttributeError and this will currently propagate instead of degrading to None as intended.
    try:
        libhip = ctypes.CDLL("libamdhip64.so")
        buf = ctypes.create_string_buffer(64)
        if libhip.hipDeviceGetPCIBusId(buf, len(buf), device_id) != 0:
            return None
        bdf = buf.value.decode()
        with open(f"/sys/bus/pci/devices/{bdf}/device") as f:
            return int(f.read().strip(), 16)
    except (OSError, ValueError, UnicodeDecodeError):
        return None

Copilot AI review requested due to automatic review settings September 1, 2026 18:26

Copilot AI 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.

Pull request overview

Copilot reviewed 13 out of 13 changed files in this pull request and generated no new comments.

Suppressed comments (3)

csrc/ck_gemm_moe_2stages_codegen/gemm_moe_tune.py:474

  • This normalization uses eval(s) on CSV cell contents. If the untuned CSV path is user-supplied (e.g., via env/arg), this is code execution; it also misses rows with leading/trailing whitespace because the startswith check doesn’t strip. Prefer parsing the known dtype strings and mapping via getattr(torch, ...) on the torch. suffix (no eval) and strip before checking.
            untunedf[col] = untunedf[col].map(
                lambda s: (
                    str(dtypes.normalize_fp8_dtype(eval(s)))
                    if isinstance(s, str) and s.startswith("torch.float8_e4m3")
                    else s

csrc/ck_gemm_a8w8/gemm_a8w8_tune.py:137

  • This normalization uses eval(s) on CSV content to obtain a torch dtype. Even if typical shape lists are trusted, untuned_gemm_file can be overridden and eval turns the tuner into a code-execution sink. Parse the known torch.float8_e4m3* strings without eval (e.g., getattr(torch, name[len('torch.'):])) and then normalize.
            untunedf["q_dtype_w"] = untunedf["q_dtype_w"].map(
                lambda s: (
                    str(dtypes.normalize_fp8_dtype(eval(s)))
                    if isinstance(s, str) and s.strip().startswith("torch.float8_e4m3")
                    else s

aiter/jit/utils/chip_info.py:462

  • _get_pci_chip_id() intends to return None when HIP/sysfs probing is unavailable, but a missing hipDeviceGetPCIBusId symbol (or other attribute lookup failure) would currently raise AttributeError and break get_device_name() dispatch on gfx942. Treat AttributeError the same as other probe failures and return None.
    except (OSError, ValueError, UnicodeDecodeError):

@zufayu
zufayu requested a review from gyohuangxin September 2, 2026 01:21
@gyohuangxin
gyohuangxin requested a review from yzhou103 September 2, 2026 02:25
Copilot AI review requested due to automatic review settings September 2, 2026 09:17

Copilot AI 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.

🟡 Changes recommended

The new unit test uses eval() on CSV-derived strings and fully materializes large CSVs into memory, creating avoidable security and performance risks in CI.

Once you've addressed the issues Copilot identified, you can request another Copilot review.

Review details

Suppressed comments (1)

Previously missed (1) — in code that hasn't changed since the last review.

op_tests/tuning_tests/test_arch_resolution.py:346

  • This test loads every tuned CSV fully into memory (rows = list(...)) just to check for cu_num=228 rows, which can make Level-1 tests slow and memory-heavy as config CSVs grow. You can stream the DictReader and use reader.fieldnames to detect the schema without materializing the whole file.
  • Files reviewed: 13/13 changed files
  • Comments generated: 1
  • Review effort level: Lite

Copilot AI review requested due to automatic review settings September 2, 2026 09:22

Copilot AI 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.

🔵 Needs a closer look

Hardware-SKU-specific dispatch/build-target logic plus large tuned-config expansions warrant final human verification on representative MI300A environments.

Review details

Suppressed comments (2)

Previously missed (2) — in code that hasn't changed since the last review.

aiter/jit/utils/build_targets.py:41

  • The comment for the gfx942 entry implies MI300A/MI308X builds always require an explicit CU_NUM override, but chip_info.get_build_targets() can now refine the CU count automatically when a live GPU is present. Clarifying that CU_NUM is required specifically for env-only/no-GPU target resolution would avoid confusion for users reading this table.
    op_tests/tuning_tests/test_arch_resolution.py:345
  • In TestTunedConfigArchConsistency, rows = list(csv.DictReader(f)) loads each tuned CSV fully into memory even though the test only needs to scan for cu_num == 228. This can be unnecessarily slow/heavy given large tuned CSVs; streaming the DictReader avoids the allocation and keeps the same behavior.
  • Files reviewed: 13/13 changed files
  • Comments generated: 0 new
  • Review effort level: Lite

@afanfa

afanfa commented Sep 14, 2026

Copy link
Copy Markdown
Contributor Author

ping

@github-actions

github-actions Bot commented Oct 7, 2026

Copy link
Copy Markdown
Contributor

No activity for 15 days, so this is now labelled stale. A push or a comment clears it; keep-open exempts it. Nothing is closed automatically.

@github-actions github-actions Bot added the stale The PR hasn't been updated in 2 weeks label Oct 7, 2026
@yzhou103

Copy link
Copy Markdown
Contributor

Hi please resolve the conflicts

Keep MI300A naming on the version-gated PciChipId lookup, include cu_num
in the OPUS parity key, and remap MI300A BF16 OPUS rows to local kernel ids.

Co-authored-by: Cursor <cursoragent@cursor.com>
Copilot AI lite review requested due to automatic review settings October 27, 2026 15:38

Copilot AI 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.

Copilot was unable to review this pull request because the user who requested the review has insufficient remaining quota.

@yzhou103

yzhou103 commented Nov 5, 2026

Copy link
Copy Markdown
Contributor

@valar-cop review

@valar-cop valar-cop left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

a few small things..

need a run on mi300a (gfx942, 228 cu) to cover the new tuned rows and the fnuz branch?

ci green.

Comment thread aiter/jit/utils/chip_info.py Resolved
@zoe-cop

zoe-cop commented Nov 5, 2026

Copy link
Copy Markdown

Follow-up review of head 44ef14392998a314433adcc93a6532a4dda5467b: two non-blocking fixes on the new content, plus a separate pre-existing CSV update issue worth tracking.

1. Reconcile the numeric IDs in six new CK-Tile rows.

In a8w8_blockscale_tuned_gemm.csv, these gfx942/228 rows save a kernelId that differs from the current ID of their kernelName:

M (N=7168, K=256) Saved ID Current ID of the saved name
1024, 1536, 4096 2 20
8192, 20480 3 21
2048 4 22

Please regenerate/reconcile these IDs against the current registry. ID-based manual reproduction would otherwise select a different candidate. The saved names exist, and codegen/runtime dispatch resolves them by name; this is stale metadata, not a demonstrated wrong-kernel runtime failure.

2. Make the FP8 regression tests assert both native encodings independently of the test host.

TestNormalizeFp8Dtype and the shipped-list check use the actual host's dtypes.fp8. The tuning-tests job uses gfx950/MI35x, so its conversion check exercises fnuz -> fn; it does not explicitly guard the fn -> fnuz direction needed by MI300A.

Please add subcases that patch the native alias to each of torch.float8_e4m3fn and torch.float8_e4m3fnuz, and assert both input flavors resolve correctly in each subcase. Keep the real-host integration check as well.

3. Separate follow-up: decimal-formatted numeric CSV keys can silently prevent an approved update (already present in base).

_merge_compare_filtered_results converts identity columns to strings on both sides. Numerically equal values such as 228/228.0 and 16/16.0 then fail to match.

Reproduction at the CSV-method level:

  • Existing tuned CSV stores integral identity fields as decimals; the untuned shape/comparison plan has integer-valued keys.
  • The candidate writer saves a new kernel. Synthetic pre/post results of 10/5 us make _build_compare_update_plan set update=True.
  • The final merge fails to select that candidate by its stringified key and retains the old kernel.

This was reproduced in head and base with native FP8 spelling, so it is not caused by the new FP8 resolver. The reproduction executes the source CSV methods with mocked device/dtype attributes and synthetic benchmark results; those times are not GPU measurements.

Please normalize integral key fields consistently across old/candidate/comparison data, validating integrality rather than truncating, and add a round-trip regression that checks the final saved kernel after write/reload. This can be a separate fix; _resolve_fp8_name does not address numeric CSV formatting.

Scoped GPU validation: on MI308X (gfx942, 80 CU), the new architecture/FP8 test file passed 25 tests with no skips. All seven CK-Tile shapes above/in the new CSV, each with two random seeds, plus M=127/129/1023/1025 boundary cases passed on both head and base (18 cases per side). Outputs were initialized to NaN and compared against a PyTorch FP32 dequantization/matmul reference converted to BF16; all elements were within atol=rtol=0.01. Head's actual shared CSV loader also converted fn -> fnuz, followed by successful GPU quantization and GEMM.

These kernel checks used a diagnostic subset build of the four saved names from unmodified kernel sources and direct native calls. They do not establish full default-registry build coverage, MI300A 228-CU automatic dispatch, performance, or a complete autotuner --update_improved run.

afanfa and others added 2 commits November 5, 2026 08:57
A multi-arch wheel built on an MI300A host replaced (gfx942, 304) with
(gfx942, 228) and lost the MI300X instances. Add the live target next to
the default instead. Also drop a stray comment in FmoeTuner.

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

# Conflicts:
#	aiter/jit/utils/build_targets.py
#	aiter/jit/utils/chip_info.py
@aiter-bot

aiter-bot commented Nov 5, 2026 •

Copy link
Copy Markdown

READY FOR HUMAN REVIEW

PR Self-Review

Revision 518476e854a6 · 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

Adds MI300A (gfx942, 228 CU) support. Before this PR, MI300A was reported as "MI300" (same as MI300X), so SKU-specific dispatch could not tell them apart. Its 228-CU tuned rows did not exist, so every GEMM/MoE shape fell back to heuristic kernels. Tuning shape lists that name `torch.float8_e4m3fn` also aborted with "Unsupported dtype" (or were silently skipped by the bpreshuffle tuner) on gfx942, which only implements e4m3fnuz.
- `chip_info`: map PCI chip id 0x74A0 to "MI300A" in `get_device_name()`.
- `dtypes.normalize_fp8_dtype()` + `TunerCommon._normalize_fp8_columns()`: resolve either e4m3 flavor in untuned shape lists to the arch-native encoding on ingest (name lookup, no `eval`), in the shared loader so every tuner family gets it.
- Tuned configs: MI300A rows for bf16, a8w8, a8w8 blockscale GEMM and fused MoE.
- Build targets: after merging main, this PR uses upstream's `GFX_CU_NUMS_MAP` (gfx942 → 304/228/80) unchanged. The earlier live-GPU CU refinement was removed, so wheels do not depend on the build host.

Validation and unverified scope

Hardware: AMD Instinct MI300A (gfx942, 228 CU), ROCm 7.14, torch 2.14.0+rocm7.14, Python 3.11.
- `python -m pytest op_tests/tuning_tests/test_arch_resolution.py op_tests/test_build_targets_architecture.py tests/test_chip_info_xcds.py` → 84 passed (this exercises the fn→fnuz direction, since the host is gfx942).
- `python -m pytest op_tests/test_opus_a16w16_policy_parity.py` → 19 passed.
- Level 0+1 suite from `tuning-tests.yaml` → all pass except 4 tests (`test_tuner_infra` GroupedFmoe x2, `test_fast_tuning` a16w16 policy, `test_mp_tuner_logic` timing) that fail identically on a clean `upstream/main` checkout in this environment (Triton version), so they are unrelated.

Unverified: the new tuned rows were not re-benchmarked at this revision; the gfx950 (fnuz→fn) direction is covered only by the CI tuning job on MI35X. The Aiter Test wheel build/GPU tests have not completed (Docker Hub 429 on the fork PR's anonymous image pull).

PR-specific question

The only CI change adds `op_tests.tuning_tests.test_arch_resolution` to the existing Level 0+1 `unittest` command in `.github/workflows/tuning-tests.yaml`. It adds no new triggers, permissions, secrets or steps. Fork PR code already runs in that same unprivileged container job, and no PR text is interpolated into a command.

Changed op tests

- `op_tests/tuning_tests/test_arch_resolution.py` (new): pins MI300A SKU naming by PCI id, the version-gated PciChipId ordinal, legacy cu_num→gfx backfill, that `GPU_ARCHS` builds include (gfx942, 228) alongside 304/80, fp8 e4m3 resolution, that every shipped shape list (incl. bpreshuffle) is tunable on the running arch, and that the cu_num=228 tuned rows resolve to gfx942. No existing test covers MI300A identity or fp8 flavor resolution of shape lists.
- `op_tests/test_opus_a16w16_policy_parity.py`: the pre-PR codegen reference map now keys on cu_num. Without it, the MI300A (228) and MI308X (80) rows collide with MI300X (304) rows for the same shape and the parity assertion compares the wrong kernel ids. The test is otherwise unchanged.
- `op_tests/tuning_tests/README.md`: documentation only; lists the new module and adds it to the suggested Level 1 command.

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.

Six gfx942/228 CK-Tile rows in a8w8_blockscale_tuned_gemm.csv stored
kernelIds that no longer match their kernelName in the current registry
(2->20, 3->21, 4->22). Dispatch resolves by name, but ID-based
reproduction would pick a different candidate.

The fp8 normalization tests only exercised the host's native encoding, so
CI on MI35X never covered fn->fnuz (gfx942). Add subcases that patch the
native alias to each e4m3 flavor; the real-host checks stay.

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

afanfa commented Nov 5, 2026

Copy link
Copy Markdown
Contributor Author

@zoe-cop Thanks. Items 1 and 2 are fixed in 67aa535. The six blockscale rows now carry the registry IDs for their kernel names (2→20, 3→21, 4→22). New subcases patch the native fp8 alias to each e4m3 flavor and assert both inputs resolve, alongside the real-host checks. Item 3 is in base, so I'll leave it to a separate fix.

@zoe-cop

zoe-cop commented Nov 5, 2026

Copy link
Copy Markdown

The six rows now store 20/20/22/20/21/21, matching the BlockPerCu 2/3/4 expansion of kernel 0 in gemm_a8w8_blockscale_cktile_instance.py, so each saved ID and kernelName agree.

The new subcases patch dtypes.fp8 to each e4m3 flavor, and normalize_fp8_dtype reads that alias at call time, so fn->fnuz is now covered on any host. Item 3 is pre-existing; deferring it is fine.

— zoe-cop · automated reply.

Comment thread aiter/configs/bf16_tuned_gemm.csv Resolved
afanfa and others added 3 commits November 9, 2026 08:54
…efactor

ROCm#5145 replaced the FlyDSL bf16 GEMM kernels with a gfx950-only
implementation. The 53 MI300A (gfx942, cu_num=228) rows that recorded
flydsl_gemm2/3/4/5_* kernels no longer parse, so lookup skipped them and
fell back to the default path; renaming them is not an option because
the new FlyDSL A16W16 backend has no gfx942 candidates.

Retuned those 53 shapes on the remaining gfx942 backends (asm, opus,
skinny, torch) with --compare --update_improved --min_improvement_pct 3,
against the production fallback the stale rows produce. 47 shapes beat
it and were confirmed in a second independent --run_config A/B; each
stale row is replaced in place by its retuned row (30 opus, 17 asm).
The other 6 shapes did not clear 3% and are removed, so they keep the
default:

    M=80/96,  N=512,  K=4096
    M=2048,   N=6144, K=2048   M=8192/16384, N=2048, K=6144
    M=8192,   N=6144, K=2048

Second-run improvement over the fallback: 6.1% to 82.8%; 30-46% for the
M<=16 decode shapes and 68-80% for most of the rest. All rows pass the
tuner's accuracy check (max err_ratio 0.031).

Triton was excluded: aiter's Triton A16W16 kernel faults on gfx942 with
Triton 3.8 / ROCm 7.14 in the tuning environment. None of the 53 rows
used Triton, and the 12 existing MI300A Triton rows are untouched.

Measured on 4x MI300A (gfx942, 228 CU), ROCm 7.14, torch 2.14.

Co-authored-by: Cursor <cursoragent@cursor.com>
Four of the six shapes left on the default in the previous commit are
large fp32-output GEMMs that ROCm#5145 made 2.2-2.7x slower than their old
FlyDSL rows. On gfx942 only one OPUS kernel supports fp32 output, so the
asm/opus/torch search could not beat the fallback for them.

hipBLASLt can. Tuned with --with-hipblaslt and the same compare gate
(>= 3% vs the production fallback), then confirmed in a second
independent --run_config A/B:

    M      N     K     out    fallback us   hipBLASLt us
    2048   6144  2048  fp32       672.9        147.8
    8192   2048  6144  fp32      3443.4        506.7
    8192   6144  2048  fp32      2791.1        501.4
    16384  2048  6144  fp32      6622.9        985.8
    96     512   4096  bf16        10.7         10.1

All four large shapes also beat their old FlyDSL records. M=80 N=512
K=4096 still has no winner and keeps the default.

hipBLASLt rows store a library solution index, which is specific to the
hipBLASLt version. These were tuned on ROCm 7.14 (4x MI300A, gfx942,
228 CU); hipb_mm raises if the index is unknown to the installed
hipBLASLt, so re-tune them if MI300A moves to a different ROCm.

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

afanfa commented Nov 10, 2026

Copy link
Copy Markdown
Contributor Author

CI Everything is green except MI35X shard 2: op_tests/test_flydsl_hgemm.py::test_hgemm_memory, case M=19 N=120 K=1088, split_k=9. Its zero-tolerance "changed-input graph" check sees 1 of 2,280 elements off by one bf16 ulp (−3.8438 vs −3.8594). The test, the check, and the case were all added in #7546 and came in with the latest merge from main. This PR doesn't touch FlyDSL or that test, and its kernel-facing changes are gfx942 tuned rows the test doesn't read. It failed on the original run and the re-run, both on linux-aiter-mi35x-1-mtck5 runners, while the same test passes on main (1eab5ae, fvzc6 runner). I can't reproduce it locally (MI300A only), so I can't tell whether it's split-K non-determinism on those runners or a real numerical difference. Could you re-run or override the gate for that job?

Comment thread aiter/configs/bf16_tuned_gemm.csv Resolved
Review feedback: hipBLASLt rows store a solution index that is specific
to the hipBLASLt build, and lookup returns them without validation, so
an index unknown to the hipBLASLt in another container raises from
getAlgosFromIndex instead of falling back. Remove the five rows added in
2c64c88; these shapes go back to the default path:

    M=96,  N=512,  K=4096
    M=2048, N=6144, K=2048   M=8192, N=6144, K=2048
    M=8192, N=2048, K=6144   M=16384, N=2048, K=6144

This reverts commit 2c64c88.

Co-authored-by: Cursor <cursoragent@cursor.com>
@zufayu
zufayu requested review from yzhou103 and removed request for gyohuangxin November 11, 2026 06:58

@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

@afanfa
afanfa merged commit e64ffc3 into ROCm:main Nov 11, 2026
64 checks passed
@afanfa
afanfa deleted the mi300a-enablement branch November 11, 2026 16:03
vpietila-amd pushed a commit that referenced this pull request Nov 12, 2026
* [HIP] Identify MI300A and resolve gfx942 CU counts correctly

gfx942 covers three SKUs with different CU counts -- MI300X (304), MI300A
(228) and MI308X (80) -- but AITER could not tell them apart, and two
separate bugs made a 228-CU part unusable with tuned configs.

_get_pci_chip_id() queried the device id by numeric ordinal:

    hipDeviceAttributePciChipId = 10019

The AMD-specific block of hipDeviceAttribute_t is an unnumbered enum, so
those ordinals shift whenever an attribute is inserted ahead of them. On
ROCm 7.2, counting the enum in hip_runtime_api.h gives 10018 =
NumberOfXccs, 10019 = MaxAvailableVgprsPerThread and 10020 = PciChipId,
so the call returned 512 on every device and the MI308 chip-id check
could never match -- MI308X has been reported as "MI300". The C++ side
uses the symbolic name and was always correct, so this is Python-only.
Resolve the id via hipDeviceGetPCIBusId plus sysfs instead, which is
stable across ROCm versions, and return None rather than raising when it
is unavailable.

get_build_targets() used GFX_CU_NUM_MAP verbatim whenever GPU_ARCHS was
set, and that table holds one canonical CU count per arch (gfx942 -> 304,
the MI300X SPX value). On a 228-CU MI300A, GPU_ARCHS=gfx942 therefore
resolved to ('gfx942', 304), gen_instances.py filtered every (gfx942,
228) row out of the tuning CSV, and those kernels were never compiled.
The failure surfaced much later as

    kernel '...' is not present in the compiled registry

on 73 of 80 blockscale shapes. Prefer the live CU count when a requested
target's arch matches the running GPU and CU_NUM was not set explicitly.
An explicit CU_NUM still wins and cross-compiles are unaffected:

    GPU_ARCHS=gfx942 (on MI300A)  [('gfx942', 304)] -> [('gfx942', 228)]
    GPU_ARCHS=gfx942 CU_NUM=304   [('gfx942', 304)] -> unchanged
    GPU_ARCHS=gfx950 (cross)      [('gfx950', 256)] -> unchanged
    GPU_ARCHS='gfx942;gfx950'                       -> [(942,228),(950,256)]

This is not MI300A-specific. build_targets.py's own comment notes that
DPX/QPX/CPX partition modes change the visible CU count, so MI300X in any
non-SPX mode hits the same cliff.

Also adds 0x74A0 as the MI300A chip id so get_device_name() returns
"MI300A", and registers 228 -> gfx942 in the legacy cu_num backfill.

Verified on 4x MI300A (gfx942:sramecc+:xnack-, 228 CU, ROCm 7.2, torch
2.13): chip id 0x74a0, get_device_name() "MI300A", get_cu_num() 228.
op_tests/tuning_tests test_tuner_infra, test_csv_validation,
test_config_shape_collision and test_compare_logic pass (68 tests,
37 subtests).

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

* [Tune] Resolve the fp8 e4m3 flavor to the arch-native encoding on ingest

Only one e4m3 encoding exists per chip: gfx942 implements e4m3fnuz (max
240, exponent bias 8) and gfx950 implements OCP e4m3fn (max 448, bias 7).
aiter/utility/dtypes.py already knows this via defaultDtypes, but the
shipped untuned shape lists name the OCP flavor unconditionally:

    aiter/configs/untuned_fmoe.csv                    22 occurrences
    aiter/configs/a8w8_untuned_gemm.csv               40 of 65 shapes
    model_configs/a8w8_blockscale_untuned_fmoe_*.csv  16-32 each

On gfx942 every such row aborts with "Unsupported dtype:
torch.float8_e4m3fn" out of _aiter_dtype_id, because _torch_to_aiter_dtype
is built from the arch-native alias and only ever contains one flavor. The
practical effect is that two thirds of the a8w8 GEMM shape list and the
entire fp8 MoE model coverage could not be tuned at all on any MI300-class
part -- which is why no MI300 rows exist for GLM-5.x, qwen3, minimax or
dsv3.

Add dtypes.normalize_fp8_dtype(), which maps either e4m3 encoding onto the
one the running chip implements. e5m2 and non-fp8 dtypes pass through
unchanged, and it declines to act when the arch is unknown (fp8 falls back
to uint8) rather than reinterpreting fp8 data as integers.

Apply it on ingest in FmoeTuner.get_untuned_gemm_list and
GemmA8W8Tuner.get_untuned_gemm_list, plus the two string->dtype conversion
sites in the MoE tuner. Normalizing on ingest rather than at point of use
matters: applied late, the tuned row is written as e4m3fnuz while the
untuned frame still says e4m3fn, so the tuner cannot match its own result
and reports the shape as never tuned.

q_dtype_w is part of the a8w8 dispatch key, so normalizing on ingest also
keeps the emitted rows and the tuned/untuned bookkeeping agreeing on one
dtype.

Verified on MI300A (gfx942, e4m3fnuz): all 40 previously-rejected a8w8
shapes now tune, and end-to-end against the default kernel path they win
between 1.10x and 2.41x.

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

* [Tune] Add MI300A (gfx942, cu_num=228) configs for GEMM and fused MoE

MI300A could not reach a single tuned row before this. AITER keys every
tuned kernel on (gfx, cu_num); MI300A reports (gfx942, 228), a pair that
appeared in no row of any tuning CSV, so every GEMM and MoE lookup missed
and fell through to the CU-blind shape heuristics. Three of the four
tables touched here had no gfx942 rows at all -- a8w8_blockscale was
6630 rows of pure gfx950, and bf16 was gfx1250 and gfx950 only.

Adds 527 measured rows:

    a8w8_blockscale_tuned_gemm.csv    78 rows    2.393x
    a8w8_tuned_gemm.csv               64 rows    1.529x
    bf16_tuned_gemm.csv              377 rows    1.367x
    model_configs/mi300a_tuned_fmoe.csv 8 rows   1.136x

Existing rows are untouched: the diff is 528 insertions and zero
deletions, and the pre-existing rows in all three modified files are
byte-identical.

Method. Every row was A/B'd end-to-end through the production op with
--run_config, never on the tuner's own us, which is measured by a
separate harness and is not comparable (at token=1 the MoE tuner reported
72us for a kernel whose real end-to-end cost was 83us against a 41us
default). A row ships only if it is numerically clean and beats the
default by at least 3% in two independent runs -- once in the A/B and
again after installation, with the module rebuilt so the tuned kernels
are actually compiled in. The second run matters: five bf16 rows cleared
the first measurement at 1.033-1.060x and came back at 0.944-1.022x on
the second, one of them 5.6% slower than the default it would have
replaced. All five sit in the same N=256/512 K=4096 family. Dropping them
raised the bf16 aggregate from 1.364x to 1.367x.

Losing rows are omitted rather than corrected. Lookup is exact-match, so
an absent shape keeps the existing default, which is the safe outcome:
51 of 433 bf16 rows, 5 of 13 MoE rows and 2 of 80 blockscale rows were
dropped this way. 14 of the bf16 drops were cases where the tuner's pick
was actually slower than the default, clustered at M = 1, 2, 4.

Largest wins: bf16 M=64 N=128 K=6144 fp32-out at 42.6 -> 8.1us (5.26x);
blockscale M=1024 N=7168 K=256 at 39.0 -> 13.9us (2.82x); a8w8 M=64
N=7168 K=8192 fp8 at 82.9 -> 34.4us (2.41x).

Not covered. Nine bf16 shapes (N=50016, K=6144, fp32 out, M=192..16384)
failed every candidate with err_ratio=1.0 under --mp 4; at M=16384 the
fp32 output alone is ~3.3 GB with four workers concurrent, so this looks
like memory pressure rather than a kernel fault. They keep the default.
Paged attention and MLA are not tuned here.

Measured on 4x MI300A (gfx942:sramecc+:xnack-, 228 CU, 6 XCC, 94 GiB
unified, ROCm 7.2, torch 2.13.0+rocm7.2), one GCD pinned for all
benchmarks. Depends on the fp8 e4m3 ingest fix in the preceding commit,
without which two thirds of the a8w8 shape list cannot be tuned on gfx942.

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

* [Tune] Add unit tests for GPU SKU resolution and fp8 flavor normalization

Covers the new functionality in the two preceding commits. Level 1, no GPU
required, so it runs in CI without MI300-class hardware: the device queries
and the live-GPU probes are mocked.

  TestDeviceName            get_device_name() maps chip ids to MI300A /
                            MI308 / MI300, leaves gfx950 and gfx1250 alone,
                            and does not crash when the id is unreadable.
  TestPciChipId             the id is read via hipDeviceGetPCIBusId plus
                            sysfs and returns None on failure. Asserts
                            hipDeviceGetAttribute is never called -- that is
                            the actual regression guard, since its ordinal
                            comes from an unnumbered enum and shifted
                            between ROCm releases.
  TestLegacyCuNumBackfill   228/304/80 -> gfx942 and 256 -> gfx950, for
                            both int and str cu_num.
  TestBuildTargets          the live CU count wins over the GFX_CU_NUM_MAP
                            default, an explicit CU_NUM still overrides it,
                            cross-compiling for a non-live arch keeps the
                            table default, multi-arch refines only the
                            matching target, and a missing GPU falls back to
                            the table. Also covers the MI308X/partition-mode
                            case, which is the same mechanism.
  TestNormalizeFp8Dtype     both e4m3 encodings resolve to the arch-native
                            one, idempotently; e5m2, bf16, fp32 and int8
                            pass through; an unknown arch passes through
                            rather than reinterpreting fp8 data as uint8.
  TestShippedShapeLists...  every fp8 shape in the shipped untuned lists
                            resolves to the native encoding, which is the
                            property that makes them tunable on gfx942.
  TestTunedConfigArch...    every tuned row keyed cu_num=228 is also keyed
                            gfx942, so the build-time codegen filter and the
                            runtime lookup cannot disagree.

23 tests, all passing on MI300A. Registered in the suite README table.

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

* [CI] Run the arch-resolution tests in the level 0+1 tuning suite

test_arch_resolution.py was added alongside the SKU-resolution and fp8
normalization changes but was not wired into CI, so it would never have
run there: the level 0+1 job names its test modules explicitly rather
than discovering them.

Adds it to the module list in tuning-tests.yaml and to the two places
that document that list, so the three stay in sync.

CI-style invocation on MI300A now reports 88 tests:

    python3 -m unittest op_tests.tuning_tests.test_csv_validation \
        op_tests.tuning_tests.test_tuner_infra \
        op_tests.tuning_tests.test_mp_tuner_logic \
        op_tests.tuning_tests.test_a6w6_tuning \
        op_tests.tuning_tests.test_arch_resolution -v

The single failure, test_mp_tuner_logic's
test_worker_records_start_only_when_task_leaves_queue, is pre-existing
and unrelated: it allows a spawn-context worker 5s to record a start
timestamp, but each spawned worker re-imports aiter, measured at 8.23s
on this machine. It reproduces unchanged at f2f8ed9.

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

* Fix Black formatting

* Addressing co-pilot recommendations

* Style fix

* [Tune] Move the MI300A fmoe rows into the generic tuned_fmoe.csv

model_configs/ is partitioned by model -- ds_v3, glm5_1, qwen3_235b,
kimik2, minimax_m3 and so on. mi300a_tuned_fmoe.csv was named for a
hardware SKU, which does not belong in that namespace.

The rows were never model-specific to begin with: they come from the
generic aiter/configs/untuned_fmoe.csv shape list and span three
unrelated geometries (6144/4096/E8, 2304/1536/E8 and 7168/256/E256/topk8).
Their tuned counterpart is the generic aiter/configs/tuned_fmoe.csv, which
already carries MI308X rows for two of those three geometries; the MI300A
rows now sit beside them, distinguished by cu_num exactly as intended.

Rewritten into that file's schema, which predates the gfx column:

  gfx         dropped; _ensure_gfx_column backfills it from cu_num via
              gfx_from_cu_num, where 228 -> gfx942 (added earlier in this
              branch). Verified the mapping is present rather than assumed.
  xbf16, flat dropped; both are 0 in all 8 rows, which is exactly what
              fused_moe assumes when the columns are absent (the reads are
              guarded by `if "xbf16" in cfg` / `if "flat" in cfg`).
  _tag        empty, matching all 617 existing rows; it only filters
              "flydsl_fallback" entries.

Checked for key collisions against the existing rows before appending:
none. The merged fmoe inventory now resolves to gfx942/cu228 x8,
gfx942/cu304 x83, gfx942/cu80 x962 and gfx950/cu256 x1417, with all 8
MI300A rows reachable.

No measurements change; this is a relocation. The 1.136x aggregate quoted
in the commit that added these rows still stands, but the path named there
is now aiter/configs/tuned_fmoe.csv.

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

* [Tune] Share the fp8 e4m3 normalization across all tuner families

Addresses two review comments.

1. The normalization was local to GemmA8W8Tuner, but
   GemmA8W8BpreShuffleTuner inherits the base CSV loader and its tuning
   paths return [] for any row whose q_dtype_w != dtypes.fp8. On gfx942
   the e4m3fn rows in the bpreshuffle shape lists were therefore skipped
   with a warning rather than resolved, leaving that GEMM family
   untunable. The scale is not marginal -- 781 rows across six shipped
   lists:

       a8w8_bpreshuffle_untuned_gemm_kimik3.csv          427
       a8w8_bpreshuffle_untuned_gemm_glm5.2.csv          142
       a8w8_bpreshuffle_untuned_gemm_glm47.csv            96
       a8w8_bpreshuffle_untuned_gemm_minimax_m3.csv       80
       a8w8_bpreshuffle_untuned_gemm_glm5_3_flash.csv     20
       a8w8_blockscale_bpreshuffle_untuned_gemm_mm2.5.csv 16

   Two other lists (glm5_a8w8_bpreshuffle, qwen38_flash_next) already
   spell e4m3fnuz by hand, which is the same problem worked around
   per-file.

   Moved into TunerCommon.get_untuned_gemm_list so every family -- a8w8,
   bpreshuffle, blockscale, batched and fmoe -- resolves on ingest, and
   dropped the two now-redundant per-tuner overrides. Verified on the
   kimik3 list: all 427 rows resolve to e4m3fnuz, 0 would be skipped.

   The resolver also no longer uses eval(). These cells come from CSV
   files whose path is caller-supplied, so eval() made the tuner a
   code-execution sink; it now looks the dtype up by name and returns
   anything unrecognised untouched. Covered by a test that feeds it
   "__import__('os').system(...)" among other junk.

2. TestTunedConfigArchConsistency required a literal gfx column, which
   tuned_fmoe.csv does not have -- it is keyed on cu_num alone and has
   its arch backfilled by gfx_from_cu_num at load time. The assertion
   failed for the 8 MI300A rows moved there in 2516405, so the Level 1
   suite could not pass. Confirmed before fixing: 527 cu_num=228 rows
   scanned, 8 failing, all in tuned_fmoe.csv.

   The test now asserts on the *resolved* arch, which is what dispatch
   actually keys on, and uses the literal column only when present. A
   companion test pins the legacy cu_num-only path explicitly and will
   fail loudly if tuned_fmoe.csv ever grows a gfx column.

Also streams the tuned CSVs instead of materializing them, as suggested
in an earlier review note, and extends the shape-list coverage to the
bpreshuffle and model_configs families via globs so newly added lists are
picked up automatically.

Three TestPciChipId cases had to change too: _load_hip_runtime() is
lru_cached, so patching ctypes.CDLL underneath it is a no-op once the
real library has been loaded. They now patch the loader itself, and the
new cases cover a missing symbol (AttributeError), a RuntimeError from
the Windows DLL search, and an absent /sys/bus/pci.

28 tests pass. CI level 0+1 on MI300A: 172 tests, the only failure the
pre-existing test_mp_tuner_logic spawn-timing case, plus 23 errors from
flydsl being absent in this harness. Black clean.

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

* [CI] Collapse nested with in test_arch_resolution to satisfy Ruff

Ruff SIM117 flagged the nested subTest / mock.patch context managers in
TestPciChipId.test_returns_none_when_hip_cannot_be_loaded. pre-checks runs
ruff with --exit-zero but pipes into reviewdog with -fail-on-error=true,
so the warning failed the job. Combine them into a single with statement.

ruff 0.16.0 (the version CI pins) is clean on every Python file in this
PR, and black is unchanged.

* [OPUS] Key the pre-PR codegen parity model on cu_num

test_shipped_tuned_selection_diff_is_exhaustive failed with
"assert 10301 == 10300" once the MI300A (gfx942, cu_num=228) rows from
this PR were shipped alongside upstream's MI308X (gfx942, cu_num=80) rows.

_pre_pr_codegen_tuned_map models the removed opus_gemm_lookup.h as a single
(arch, M, N, K, outdtype) table built from every shipped row, last write
wins. That is not what the pre-PR codegen did: gen_instances.py filtered
the tuned CSV to the host's cu_num before emitting the table,

    if cu_num is not None:
        tune_df = tune_df[tune_df["cu_num"] == cu_num].reset_index()

so each build saw exactly one cu_num per arch. Merging both gfx942 SKUs
into one table made the test compare MI300A's tuned kernel against the one
MI308X picked for the same shape -- dsv4_bf16_tuned_gemm.csv sorts after
bf16_tuned_gemm.csv, so its cu_num=80 row overwrote the cu_num=228 one.

Five shipped shapes are affected, all gfx942, each with a distinct
measured winner per SKU:

    M    N     K     cu_num=80 (MI308X)  cu_num=228 (MI300A)
    1    2048  4096  10301               10300
    2    2048  4096  10301               10300
    4    2048  4096  10301               10300
    8    2048  4096  10301               10300
    256  1024  4096  10210               10201

Both selections are correct for their hardware. Rewriting the MI300A rows
to MI308X's kids would make the test pass by shipping a kernel that was
never measured on MI300A, overriding the one that won the A/B there.

Key the codegen model on cu_num, matching the per-host filter and the
companion _pre_pr_python_tuned_map, which already includes cu_num. Also
make the assertion name the key and both kids on failure.

Reproduced the CI failure against the previous version of the test
(1 failed, 18 passed); all 19 pass now. ruff 0.16.0 and black clean.

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

* Keep table-default CU target when adding the live GPU's CU count

A multi-arch wheel built on an MI300A host replaced (gfx942, 304) with
(gfx942, 228) and lost the MI300X instances. Add the live target next to
the default instead. Also drop a stray comment in FmoeTuner.

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

* Reconcile MI300A blockscale kernel IDs; pin both fp8 directions in tests

Six gfx942/228 CK-Tile rows in a8w8_blockscale_tuned_gemm.csv stored
kernelIds that no longer match their kernelName in the current registry
(2->20, 3->21, 4->22). Dispatch resolves by name, but ID-based
reproduction would pick a different candidate.

The fp8 normalization tests only exercised the host's native encoding, so
CI on MI35X never covered fn->fnuz (gfx942). Add subcases that patch the
native alias to each e4m3 flavor; the real-host checks stay.

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

* [Tune] Retune MI300A bf16 GEMM rows left stale by the FlyDSL A16W16 refactor

#5145 replaced the FlyDSL bf16 GEMM kernels with a gfx950-only
implementation. The 53 MI300A (gfx942, cu_num=228) rows that recorded
flydsl_gemm2/3/4/5_* kernels no longer parse, so lookup skipped them and
fell back to the default path; renaming them is not an option because
the new FlyDSL A16W16 backend has no gfx942 candidates.

Retuned those 53 shapes on the remaining gfx942 backends (asm, opus,
skinny, torch) with --compare --update_improved --min_improvement_pct 3,
against the production fallback the stale rows produce. 47 shapes beat
it and were confirmed in a second independent --run_config A/B; each
stale row is replaced in place by its retuned row (30 opus, 17 asm).
The other 6 shapes did not clear 3% and are removed, so they keep the
default:

    M=80/96,  N=512,  K=4096
    M=2048,   N=6144, K=2048   M=8192/16384, N=2048, K=6144
    M=8192,   N=6144, K=2048

Second-run improvement over the fallback: 6.1% to 82.8%; 30-46% for the
M<=16 decode shapes and 68-80% for most of the rest. All rows pass the
tuner's accuracy check (max err_ratio 0.031).

Triton was excluded: aiter's Triton A16W16 kernel faults on gfx942 with
Triton 3.8 / ROCm 7.14 in the tuning environment. None of the 53 rows
used Triton, and the 12 existing MI300A Triton rows are untouched.

Measured on 4x MI300A (gfx942, 228 CU), ROCm 7.14, torch 2.14.

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

* [Tune] Recover MI300A large fp32-out bf16 GEMMs with hipBLASLt

Four of the six shapes left on the default in the previous commit are
large fp32-output GEMMs that #5145 made 2.2-2.7x slower than their old
FlyDSL rows. On gfx942 only one OPUS kernel supports fp32 output, so the
asm/opus/torch search could not beat the fallback for them.

hipBLASLt can. Tuned with --with-hipblaslt and the same compare gate
(>= 3% vs the production fallback), then confirmed in a second
independent --run_config A/B:

    M      N     K     out    fallback us   hipBLASLt us
    2048   6144  2048  fp32       672.9        147.8
    8192   2048  6144  fp32      3443.4        506.7
    8192   6144  2048  fp32      2791.1        501.4
    16384  2048  6144  fp32      6622.9        985.8
    96     512   4096  bf16        10.7         10.1

All four large shapes also beat their old FlyDSL records. M=80 N=512
K=4096 still has no winner and keeps the default.

hipBLASLt rows store a library solution index, which is specific to the
hipBLASLt version. These were tuned on ROCm 7.14 (4x MI300A, gfx942,
228 CU); hipb_mm raises if the index is unknown to the installed
hipBLASLt, so re-tune them if MI300A moves to a different ROCm.

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

* [Tune] Drop the MI300A hipBLASLt bf16 rows

Review feedback: hipBLASLt rows store a solution index that is specific
to the hipBLASLt build, and lookup returns them without validation, so
an index unknown to the hipBLASLt in another container raises from
getAlgosFromIndex instead of falling back. Remove the five rows added in
2c64c88; these shapes go back to the default path:

    M=96,  N=512,  K=4096
    M=2048, N=6144, K=2048   M=8192, N=6144, K=2048
    M=8192, N=2048, K=6144   M=16384, N=2048, K=6144

This reverts commit 2c64c88.

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

---------

Co-authored-by: Cursor <cursoragent@cursor.com>
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.

6 participants