Skip to content

[Triton/Gluon] [Conv2D] Optimize conv2d for RDNA - kernels, routing, and input packing - #4869

Merged
saeid-rostami merged 10 commits into
ROCm:mainfrom
saeid-rostami:conv2d-rdna-optimization-v1
Aug 26, 2026
Merged

saeid-rostami merged 10 commits into
ROCm:mainfrom
saeid-rostami:conv2d-rdna-optimization-v1

Conversation

@saeid-rostami

Copy link
Copy Markdown
Contributor

Motivation

Optimization of conv2d for RDNA architectures.

Technical Details

Summary

This PR improves the Triton Conv2D path on AMD RDNA GPUs for FP16/BF16 and
NCHW/NHWC workloads. It adds a repack-free NCHW 3x3 kernel, replaces the
multi-operation NCHW-to-NCHWc conversion with a single Triton kernel, improves
shape/layout/dtype-aware configuration selection, and adds tuned configurations
for gfx1100, gfx1151, and gfx1201.

The public API is unchanged.

What changed

  • Add a direct NCHW 3x3 kernel that reads contiguous NCHW activations without
    materializing an NCHWc input buffer.
  • Route direct NCHW only where it is profitable:
    • gfx1100 and gfx1151 use exact-shape route pins.
    • gfx1201 uses the configured spatial crossover.
    • Other shapes retain the cblocked fallback.
  • Fuse NCHW padding and layout conversion into one Triton
    _nchw_to_cblocked_kernel launch, with per-architecture CONV-PREPACK JSON
    configurations.
  • Avoid unnecessary hot-path work by allocating channels-last outputs directly,
    loading bias in its original dtype and promoting it inside the Triton
    epilogue, and skipping logging string construction when INFO logging is off.
  • Strengthen weight-pack cache keys with device, shape, stride, dtype, tensor
    version, and block size so aliases and in-place updates cannot produce stale
    cache hits.
  • Extend Conv2D config lookup with layout/dtype-specific exact-shape variants,
    followed by generic shape, M_LEQ_*, and any fallbacks.
  • Keep production launches deterministic by removing runtime autotune decorators
    and candidate tables from the kernel modules. Production always consumes the
    frozen per-architecture JSON configurations.
  • Add regression coverage for exact-route selection, fallback routing, complete
    shape keys, and config-variant precedence.

Performance

Tested on ResNet-50, Stable Diffusion 3.5 Medium, and
FLUX.2 Klein 9B across both FP16 and BF16.

  • Each model/layout/dtype combination ran in a fresh process with a fresh
    Triton disk cache; JIT compilation is excluded from the reported kernel time.
  • The benchmark flushes L2 between timed measurements.
  • NCHW optimized totals use kernel+repack, so the NCHWc input-pack cost is
    included. NHWC totals are kernel-only because that path has no separate input
    pack.

gfx1201

Layout Main total time (ms) PR total time (ms) Time reduction Main aggregate TFLOPS PR aggregate TFLOPS TFLOPS gain
NCHW 77.45 64.18 17.1% 70.52 85.11 20.7%
NHWC 92.45 52.59 43.1% 59.08 103.86 75.8%

gfx1151

Layout Main total time (ms) PR total time (ms) Time reduction Main aggregate TFLOPS PR aggregate TFLOPS TFLOPS gain
NCHW 294.989 269.364 8.69% 18.52 20.28 9.51%
NHWC 343.757 247.980 27.86% 15.89 22.03 38.62%

gfx1100

Layout Main total time (ms) PR total time (ms) Time reduction Main aggregate TFLOPS PR aggregate TFLOPS TFLOPS gain
NCHW 83.66 79.06 5.5% 65.29 69.09 5.8%
NHWC 90.86 73.39 19.2% 60.12 74.43 23.8%

Test Plan

Test Result

Submission Checklist

@saeid-rostami
saeid-rostami requested review from a team and a lite review from Copilot August 20, 2026 00:04
@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; main branch always runs both MI35X and MI300X
ci:sglang SGLang integration tests: DeepSeek-R1-MXFP4 accuracy, Qwen 3.5 accuracy
ci:atom ATOM benchmark: DeepSeek-R1-0528, GPT-OSS-120B
ci:atom_full ATOM accuracy suite for PR and main models from ATOM models_accuracy.json
ci:vllm vLLM benchmark: GPT-OSS-120B, DeepSeek-R1-0528, Kimi-K2.5
ci:all All standard extended tests (excludes ci:atom_full)

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

PR title tags:
Component tags ([Triton/Gluon], [HIP], [CK], [ASM], ...) are added to the PR title automatically from the changed files and re-synced on every push — change-type tags like [fix]/[Perf] and op tags like [MLA] are left untouched. Add the no-auto-title label to opt this PR out of title tagging.

@github-actions github-actions Bot changed the title [Triton][Conv2D] Optimize conv2d for RDNA - kernels, routing, and input packing [Triton/Gluon] [Conv2D] Optimize conv2d for RDNA - kernels, routing, and input packing Aug 20, 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

This PR optimizes the Triton Conv2D implementation for AMD RDNA architectures by adding a repack-free direct NCHW 3×3 kernel, replacing the Python-side NCHW→NCHWc conversion with a single Triton packing kernel, and strengthening configuration/routing selection (including new per-arch JSON configs and routing regression tests) while keeping the public API unchanged.

Changes:

  • Add a direct NCHW 3×3 route and kernel, and update routing to select it only when profitable (exact-shape pins or spatial crossover).
  • Fuse NCHW→NCHWc activation packing into one Triton kernel and add CONV-PREPACK per-arch configs.
  • Extend conv config lookup with layout/dtype variants and add routing/config regression tests + bench labeling updates.

Reviewed changes

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

Show a summary per file
File Description
op_tests/triton_tests/conv/test_conv2d.py Adds routing/config-lookup regression tests and updates suite documentation/counts.
op_tests/triton_tests/conv/_helpers.py Registers new direct method and adds a guard tied to config availability.
op_tests/op_benchmarks/triton/bench_conv2d.py Updates route labeling helper to pass padding into routing.
aiter/ops/triton/utils/conv_config_utils.py Adds variant-aware config lookup, prepack shape-key formatting, and config presence probing.
aiter/ops/triton/conv/README.md Updates user-facing docs to reflect new kernel family, routing, and config tiers.
aiter/ops/triton/conv/DESIGN.md Updates design documentation for direct NCHW routing, fused prepack, and config lookup behavior.
aiter/ops/triton/conv/conv2d.py Adds direct NCHW 3×3 routing and wrapper, refactors 3×3 selection logic, and reduces logging overhead.
aiter/ops/triton/conv/_utils.py Allocates NHWC outputs directly as channels-last; changes bias handling to avoid eager dtype conversion.
aiter/ops/triton/conv/_prepack.py Strengthens weight-pack cache keys and switches NCHW→NCHWc packing to a Triton kernel launch.
aiter/ops/triton/conv/_launch.py Adds fused NCHW→NCHWc pack launcher, adds direct NCHW 3×3 launcher, and passes config variants for selection.
aiter/ops/triton/configs/conv/gfx950-CONV-PREPACK.json Adds tuned prepack configs for gfx950.
aiter/ops/triton/configs/conv/gfx942-CONV-PREPACK.json Adds tuned prepack configs for gfx942.
aiter/ops/triton/configs/conv/gfx1250-CONV-PREPACK.json Adds tuned prepack configs for gfx1250.
aiter/ops/triton/configs/conv/gfx1201-CONV-PREPACK.json Adds tuned prepack configs for gfx1201.
aiter/ops/triton/configs/conv/gfx1200-CONV-PREPACK.json Adds tuned prepack configs for gfx1200.
aiter/ops/triton/configs/conv/gfx1151-CONV-PREPACK.json Adds tuned prepack configs for gfx1151.
aiter/ops/triton/configs/conv/gfx1100-CONV-PREPACK.json Adds tuned prepack configs for gfx1100.
aiter/ops/triton/configs/conv/gfx1201-CONV-GENERAL.json Updates tuned parameters (e.g., BLOCK_K) for gfx1201 general kernel.
aiter/ops/triton/configs/conv/gfx1151-CONV-GENERAL.json Updates tuned parameters (e.g., BLOCK_K) for gfx1151 general kernel.
aiter/ops/triton/configs/conv/gfx1201-CONV-3X3-NHWC.json Adds dtype-specific exact-shape pins for gfx1201 NHWC 3×3 config.
aiter/ops/triton/configs/conv/gfx1201-CONV-3X3-NCHW.json Adds new gfx1201 direct NCHW 3×3 config table.
aiter/ops/triton/configs/conv/gfx1201-CONV-3X3-CBLOCKED.json Adds fp16 variant pins and updates tuned parameters for gfx1201 cblocked 3×3 config.
aiter/ops/triton/configs/conv/gfx1201-CONV-1X1.json Adds NHWC exact-shape pins for gfx1201 1×1 config.
aiter/ops/triton/configs/conv/gfx1151-CONV-WINO-F4X3-OUTPUT.json Updates tuned Winograd output-transform config for gfx1151.
aiter/ops/triton/configs/conv/gfx1151-CONV-WINO-F4X3-INPUT.json Updates tuned Winograd input-transform config for gfx1151.
aiter/ops/triton/configs/conv/gfx1151-CONV-WINO-F4X3-GEMM.json Updates tuned Winograd GEMM config for gfx1151.
aiter/ops/triton/configs/conv/gfx1151-CONV-3X3-NHWC.json Updates tuned NHWC 3×3 config for gfx1151.
aiter/ops/triton/configs/conv/gfx1151-CONV-3X3-NCHW.json Adds new gfx1151 direct NCHW 3×3 exact-route config table.
aiter/ops/triton/configs/conv/gfx1151-CONV-3X3-CBLOCKED.json Updates tuned cblocked 3×3 config for gfx1151.
aiter/ops/triton/configs/conv/gfx1151-CONV-1X1.json Updates tuned 1×1 config for gfx1151 and adds variant shape maps.
aiter/ops/triton/configs/conv/gfx1100-CONV-3X3-NCHW.json Adds new gfx1100 direct NCHW 3×3 exact-route config table.
aiter/ops/triton/_triton_kernels/conv/nchw_to_cblocked.py Adds new Triton kernel for fused NCHW→NCHWc activation packing.
aiter/ops/triton/_triton_kernels/conv/helpers.py Removes runtime autotune gate module from production Conv2D kernels.
aiter/ops/triton/_triton_kernels/conv/conv_general.py Removes runtime autotune paths, adds variants plumbing, and promotes bias to fp32 in-kernel.
aiter/ops/triton/_triton_kernels/conv/conv_3x3.py Removes runtime autotune paths, adds direct NCHW 3×3 kernel, variants plumbing, and promotes bias to fp32 in-kernel.
aiter/ops/triton/_triton_kernels/conv/conv_3x3_winograd_f4x3.py Removes runtime autotune paths and promotes bias to fp32 in output transform.
aiter/ops/triton/_triton_kernels/conv/conv_1x1.py Removes runtime autotune paths, adds variants plumbing, and promotes bias to fp32 in-kernel.

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

@zufayu
zufayu requested a review from Dewei-Wang-sh August 20, 2026 00:52
Copilot AI review requested due to automatic review settings August 20, 2026 01:19

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 37 out of 37 changed files in this pull request and generated no new comments.

Suppressed comments (2)

aiter/ops/triton/_triton_kernels/conv/nchw_to_cblocked.py:16

  • The new Triton kernel _nchw_to_cblocked_kernel is declared with @triton.jit but does not set a config-aware repr (e.g. via make_kernel_repr(...)) like the other conv kernels. Without a stable repr embedding the tuned meta-parameters, trace/profiling labels can be ambiguous and compiled-kernel caching/debuggability can regress compared to existing kernels in this directory.
def _get_config(shape_key=None, M=None, variants=()):
    return get_conv_config("CONV-PREPACK", shape_key=shape_key, M=M, variants=variants)


@triton.jit
def _nchw_to_cblocked_kernel(
    X,

aiter/ops/triton/conv/conv2d.py:620

  • conv2d_nchw_3x3_direct() can be called directly, but it unconditionally performs config lookup for CONV-3X3-NCHW via _launch_3x3_nchw even on architectures where that optional JSON table is not shipped. On such devices this will fail with a low-level config-file load error rather than a clear, actionable exception. Consider guarding with has_conv_config("CONV-3X3-NCHW") and raising a ValueError explaining that direct NCHW 3x3 is only available on specific archs/configs (and that the caller should use conv2d_nchw routing instead).

@Ragua1

Ragua1 commented Aug 21, 2026 •

Copy link
Copy Markdown
Contributor

Measured on gfx1150 (Radeon 890M, RDNA 3.5), Windows with native ROCm, torch 2.11.0+rocm7.15.0a20260728, Triton 3.7.1, on 2026-08-21 against main at ba59a37 and this branch at e363cf6. I maintain a downstream consumer that routes diffusion VAE-decode convolutions through aiter.ops.triton.conv. My primary interest is gfx1101 (RX 7800 XT), which I cannot measure at the moment; gfx1150 is the card I have to hand, and neither ships a conv table, so they are in the same position with respect to this PR.

The performance direction of this PR is good on this part, which is why I would like it to stay reachable. On the routes it keeps — cblocked, fp16, N=1, K=C, interleaved A-B-A-B in one session, the two runs per revision shown rather than averaged:

shape main this PR change
C=512, 128×128, 3×3 10.575 / 10.640 ms 10.520 / 10.482 ms −1 %
C=512, 256×256, 3×3 41.706 / 41.043 ms 40.992 / 40.567 ms −1.5 %
C=256, 512×512, 3×3 41.651 / 42.563 ms 39.344 / 39.325 ms −6.6 %
C=128, 1024×1024, 3×3 43.978 / 44.401 ms 37.485 / 37.546 ms −15.0 %

The new packing kernel accounts for it: pack cost drops from 9.37 to 5.85 ms at 1024×1024, about 38 % cheaper than the host-side permute().contiguous() on main. Separately, conv2d_1x1(layout="nhwc") is 2.3× faster here (2.85 → 1.25 ms).

The question

On an architecture that ships no tuned table, the two revisions behave differently:

main @ ba59a37,    no config, AITER_TRITON_CONV_AUTOTUNE=0  ->  FileNotFoundError
main @ ba59a37,    no config, AITER_TRITON_CONV_AUTOTUNE=1  ->  runs, rel_l2 7.67e-04
this PR @ e363cf6, no config, AITER_TRITON_CONV_AUTOTUNE=1  ->  FileNotFoundError

Two changes combine: _triton_kernels/conv/helpers.py is deleted along with the seven AUTOTUNE_*_CONFIGS, so the env-var bypass documented in conv/README.md no longer exists; and CONV-PREPACK is a newly required per-arch family on a path that previously read no JSON at all, so it is the first lookup to fail. The per-architecture requirement goes from 7 tables to 9, and every shipped table is arch-specific. The updated README states which architectures ship one — gfx1100, gfx1151, gfx1200 and gfx1201 — and that is four of the eight in RDNA_ARCHS; the other four are gfx1030, gfx1101, gfx1102 and gfx1150, so conv2d becomes unreachable on RX 6800/6900, RX 7800 XT and 7700 XT, RX 7600 and Strix Point.

I am not claiming any of those is documented as supported; op_tests/triton_tests/conv/_helpers.py lists exactly the four that have tables, so the current state is self-consistent. The question is what the intended path is for the other four now that retuning moves out of the package, since the README paragraph describing how to run without a table is removed in this diff, and its replacement states that "candidate search spaces and the tuning driver live outside the production package".

Two observations that may make this cheap to resolve. First, on this part the config contents matter much less than their presence: a 48-combination sweep over BLOCK_M/BLOCK_N/BLOCK_K/GROUP_SIZE_M/num_warps on the four shapes above — route asserted as DIRECT_NCHW_3X3, output checked against torch.nn.functional.conv2d — reaches parity with the cblocked route at best (0.93–1.16× over three repeats), and simply borrowing a shipped table gives much the same (gfx1151 0.84×, gfx1201 0.85×). Second, it would have to be a curated default rather than arbitrary values: in the same sweep num_warps=2 with BLOCK_N=128 produced 0.05×, twenty times slower than cblocked. A nearest-architecture fallback would therefore cost very little and remove a hard failure — and has_conv_config(), which this PR adds with load_config_json(..., required=False) and a docstring calling the table optional, is already the primitive for it. (The suppressed automated comment on conv2d.py proposes exactly that as a guard for conv2d_nchw_3x3_direct(); that is one entry point, whereas on an architecture with no table at all the first mandatory get_conv_config fails inside load_config_json before any KeyError path is reached.)

On gfx1101 the picture came out a little differently — there the table contents did move the number. Measured in early August 2026 on the pre-#4869 tree, not re-verified since, and to one significant figure: repeating the comparison in a later session shifted the ratio by about two points, so I would not quote a second digit. The shipped gfx1100 table from #4265 was indistinguishable from an RDNA4 table on that card — a ratio of 1.00 against a 6–8 % spread within a single arm — so same-generation donorship was not the lever. Changing one value in the RDNA4 table — BLOCK_K 64→32 in CONV-3X3-CBLOCKED — was worth about 10 % there, while the same edit applied to the gfx1100 table was worth nothing, so the useful value is donor-specific rather than portable. And a held-out neighbour table cost about 9 % against one tuned on the card. So a neighbour table is worth far more than a hard failure, and still leaves something on the table; that ordering is what matters here.

Is a curated nearest-architecture fallback something you would consider, or is the intent that an architecture without a table goes through the offline tuning path first? Either answer is useful — it tells downstream consumers on those four architectures what to plan for.

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

Sharing some comments.

Comment thread aiter/ops/triton/conv/_utils.py Resolved
Comment thread aiter/ops/triton/conv/_launch.py Resolved
Comment thread aiter/ops/triton/conv/conv2d.py Resolved
- share tuning configurations between FP16 and BF16
- remove dtype-specific variant dispatch while preserving layout overrides
- use the common activation helper across Conv2D kernels
- simplify output allocation for NCHW and NHWC layouts
- switch Conv2D logging to lazy placeholder formatting
- update configuration tests and Conv2D documentation
Copilot AI review requested due to automatic review settings August 26, 2026 17:08
@saeid-rostami

Copy link
Copy Markdown
Contributor Author

Measured on gfx1150 (Radeon 890M, RDNA 3.5), Windows with native ROCm, torch 2.11.0+rocm7.15.0a20260728, Triton 3.7.1, on 2026-08-21 against main at ba59a37 and this branch at e363cf6. I maintain a downstream consumer that routes diffusion VAE-decode convolutions through aiter.ops.triton.conv. My primary interest is gfx1101 (RX 7800 XT), which I cannot measure at the moment; gfx1150 is the card I have to hand, and neither ships a conv table, so they are in the same position with respect to this PR.

The performance direction of this PR is good on this part, which is why I would like it to stay reachable. On the routes it keeps — cblocked, fp16, N=1, K=C, interleaved A-B-A-B in one session, the two runs per revision shown rather than averaged:

shape main this PR change
C=512, 128×128, 3×3 10.575 / 10.640 ms 10.520 / 10.482 ms −1 %
C=512, 256×256, 3×3 41.706 / 41.043 ms 40.992 / 40.567 ms −1.5 %
C=256, 512×512, 3×3 41.651 / 42.563 ms 39.344 / 39.325 ms −6.6 %
C=128, 1024×1024, 3×3 43.978 / 44.401 ms 37.485 / 37.546 ms −15.0 %
The new packing kernel accounts for it: pack cost drops from 9.37 to 5.85 ms at 1024×1024, about 38 % cheaper than the host-side permute().contiguous() on main. Separately, conv2d_1x1(layout="nhwc") is 2.3× faster here (2.85 → 1.25 ms).

The question

On an architecture that ships no tuned table, the two revisions behave differently:

main @ ba59a37,    no config, AITER_TRITON_CONV_AUTOTUNE=0  ->  FileNotFoundError
main @ ba59a37,    no config, AITER_TRITON_CONV_AUTOTUNE=1  ->  runs, rel_l2 7.67e-04
this PR @ e363cf6, no config, AITER_TRITON_CONV_AUTOTUNE=1  ->  FileNotFoundError

Two changes combine: _triton_kernels/conv/helpers.py is deleted along with the seven AUTOTUNE_*_CONFIGS, so the env-var bypass documented in conv/README.md no longer exists; and CONV-PREPACK is a newly required per-arch family on a path that previously read no JSON at all, so it is the first lookup to fail. The per-architecture requirement goes from 7 tables to 9, and every shipped table is arch-specific. The updated README states which architectures ship one — gfx1100, gfx1151, gfx1200 and gfx1201 — and that is four of the eight in RDNA_ARCHS; the other four are gfx1030, gfx1101, gfx1102 and gfx1150, so conv2d becomes unreachable on RX 6800/6900, RX 7800 XT and 7700 XT, RX 7600 and Strix Point.

I am not claiming any of those is documented as supported; op_tests/triton_tests/conv/_helpers.py lists exactly the four that have tables, so the current state is self-consistent. The question is what the intended path is for the other four now that retuning moves out of the package, since the README paragraph describing how to run without a table is removed in this diff, and its replacement states that "candidate search spaces and the tuning driver live outside the production package".

Two observations that may make this cheap to resolve. First, on this part the config contents matter much less than their presence: a 48-combination sweep over BLOCK_M/BLOCK_N/BLOCK_K/GROUP_SIZE_M/num_warps on the four shapes above — route asserted as DIRECT_NCHW_3X3, output checked against torch.nn.functional.conv2d — reaches parity with the cblocked route at best (0.93–1.16× over three repeats), and simply borrowing a shipped table gives much the same (gfx1151 0.84×, gfx1201 0.85×). Second, it would have to be a curated default rather than arbitrary values: in the same sweep num_warps=2 with BLOCK_N=128 produced 0.05×, twenty times slower than cblocked. A nearest-architecture fallback would therefore cost very little and remove a hard failure — and has_conv_config(), which this PR adds with load_config_json(..., required=False) and a docstring calling the table optional, is already the primitive for it. (The suppressed automated comment on conv2d.py proposes exactly that as a guard for conv2d_nchw_3x3_direct(); that is one entry point, whereas on an architecture with no table at all the first mandatory get_conv_config fails inside load_config_json before any KeyError path is reached.)

On gfx1101 the picture came out a little differently — there the table contents did move the number. Measured in early August 2026 on the pre-#4869 tree, not re-verified since, and to one significant figure: repeating the comparison in a later session shifted the ratio by about two points, so I would not quote a second digit. The shipped gfx1100 table from #4265 was indistinguishable from an RDNA4 table on that card — a ratio of 1.00 against a 6–8 % spread within a single arm — so same-generation donorship was not the lever. Changing one value in the RDNA4 table — BLOCK_K 64→32 in CONV-3X3-CBLOCKED — was worth about 10 % there, while the same edit applied to the gfx1100 table was worth nothing, so the useful value is donor-specific rather than portable. And a held-out neighbour table cost about 9 % against one tuned on the card. So a neighbour table is worth far more than a hard failure, and still leaves something on the table; that ordering is what matters here.

Is a curated nearest-architecture fallback something you would consider, or is the intent that an architecture without a table goes through the offline tuning path first? Either answer is useful — it tells downstream consumers on those four architectures what to plan for.

@Ragua1
Thanks for the detailed measurements and analysis. It is especially useful to see that the packing kernel improves the larger cblocked workloads on gfx1150.
The removal of AITER_TRITON_CONV_AUTOTUNE and the in-kernel autotuning configurations was intentional. After discussing this internally, I was asked to follow the AITER policy that production Triton kernels should not perform runtime autotuning or carry autotuning search spaces.
The intended workflow for an architecture without Conv2D configuration files is to tune the kernels offline, add the resulting per-architecture JSON configuration files, and submit them to AITER in a PR. We do not currently intend to select an arbitrary or nearest-architecture configuration automatically, since—as your gfx1101 results also demonstrate—configurations are not always portable and a poorly matched fallback can cause a substantial performance regression.
Therefore, the current failure on architectures without the required configuration tables is expected under this policy. If you would like to add support for gfx1101, gfx1102, gfx1150, or another missing architecture, please tune and submit the corresponding Conv2D configuration files and add me as a reviewer. I would be happy to help review and validate them.

Copilot stopped reviewing on behalf of saeid-rostami due to an error August 26, 2026 17:35

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 encountered an error and was unable to review this pull request. You can try again by re-requesting a review.

Copilot AI review requested due to automatic review settings August 26, 2026 19:21

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 37 out of 37 changed files in this pull request and generated 1 comment.

Suppressed comments (1)

aiter/ops/triton/_triton_kernels/conv/nchw_to_cblocked.py:15

  • New Triton kernel _nchw_to_cblocked_kernel is missing a config-aware repr=make_kernel_repr(...). Other Conv2D kernels set a repr so traces embed the tuned config; without it, profiling/debugging can’t reliably distinguish different BLOCK_C/BLOCK_M configs for this kernel.
@triton.jit
def _nchw_to_cblocked_kernel(

Comment thread aiter/ops/triton/conv/conv2d.py Resolved

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 37 out of 37 changed files in this pull request and generated no new comments.

Suppressed comments (1)

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

aiter/ops/triton/conv/conv2d.py:160

  • In _resolve_route's NHWC branch, the router ignores the "general" result from _select_3x3_method (which returns "general" for narrow C). This makes the narrow-channel heuristic ineffective for NHWC 3x3 and can route small-C workloads to the NHWC 3x3 kernel even when the selector chose the general kernel.

@brunomazzottiamd brunomazzottiamd 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!

Copilot AI review requested due to automatic review settings August 26, 2026 20:34

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 37 out of 37 changed files in this pull request and generated 1 comment.

Comment thread aiter/ops/triton/conv/_launch.py Resolved
@saeid-rostami
saeid-rostami merged commit 2f9ee5e into ROCm:main Aug 26, 2026
76 checks passed
Boss2002n added a commit that referenced this pull request Aug 27, 2026
Move all 59 flat configs/conv/<arch>-<CONFIG_NAME>.json files to
configs/<arch>/triton/conv/<d_type>/DEFAULT.json, the layout GEMM already
uses, and point _conv_config_path() at the shared resolve_config_dir()
probe instead of building the legacy path by hand. This picks up the ten
tables #4869 added (CONV-PREPACK on all seven arches, CONV-3X3-NCHW on
gfx1100/gfx1151/gfx1201) alongside the original 49.

The renames and the loader flip land in one commit so every revision
resolves conv configs from exactly one layout: no legacy_dir fallback is
needed and bisect stays clean. File contents are untouched (pure renames).
_conv_config_path() is the single choke point, so get_conv_config(),
has_conv_config(), conv_config_uses_exact_routes() and
has_exact_conv_config() all pick up the nested path; the variant-aware
four-tier walk, STANDARD_M_BOUNDS and the lru_caches are untouched.
brunomazzottiamd pushed a commit that referenced this pull request Aug 27, 2026
Move all 59 flat configs/conv/<arch>-<CONFIG_NAME>.json files to
configs/<arch>/triton/conv/<d_type>/DEFAULT.json, the layout GEMM already
uses, and point _conv_config_path() at the shared resolve_config_dir()
probe instead of building the legacy path by hand. This picks up the ten
tables #4869 added (CONV-PREPACK on all seven arches, CONV-3X3-NCHW on
gfx1100/gfx1151/gfx1201) alongside the original 49.

The renames and the loader flip land in one commit so every revision
resolves conv configs from exactly one layout: no legacy_dir fallback is
needed and bisect stays clean. File contents are untouched (pure renames).
_conv_config_path() is the single choke point, so get_conv_config(),
has_conv_config(), conv_config_uses_exact_routes() and
has_exact_conv_config() all pick up the nested path; the variant-aware
four-tier walk, STANDARD_M_BOUNDS and the lru_caches are untouched.
Boss2002n added a commit that referenced this pull request Aug 27, 2026
utils/ reorganized for reuse: ONE config_utils.py absorbs core.py and the
gemm/conv/mhc/tuned config_utils modules, laid out in five clearly-banded
sections (shared core, GEMM, conv, MHC, tuned kernel entries) so each op
family stays visibly separate inside the single import surface. The
duplicate utils/mha_kernel_utils.py is dropped in favor of the byte-equal
utils/_triton copy; every importer tree-wide is updated. gmm_common stays
put under utils/ so tests and benchmarks keep importing a public module.
The MHC section is absorbed verbatim and still reads the legacy flat
layout -- the MHC config-file migration later flips only this section's
path builders.

resolve_config_dir() gets the contract the #4947/#5019 reviews asked for,
hardened: an Args/Returns/Raises/Example docstring; every argument
validated against a whitelist with assertions (op, config_name and the
programmer-written arch= override must match strict identifier patterns,
the driver-derived running arch tolerates vendor formats but must stay
path-safe, backend is checked against ("triton", "gluon")) so a bad value
fails closed instead of resolving an escaped or wrong directory; a single
str return value (the unused name_prefix tuple element is gone, no caller
unpacks a tuple); and fully deterministic behavior -- the directory is
built, never probed, so nothing silently returns a bogus last candidate:
a missing required table fails loudly in load_config_json() naming the
exact nested path, and optional tables pass required=False.

The MOE dispatch loaders now declare their backend per dispatch path
(a8w4's triton path resolves backend="triton", the gluon paths "gluon";
a4w4 is gluon-only), and gfx1250 gains byte-identical triton seeds copied
from gfx950 (GEMM-AFP4WFP4, moe A8W4, HSTU fwd/bwd) so the one resolution
the caller-declared policy changes -- triton GEMM-AFP4WFP4 on gfx1250 no
longer borrowing gluon tables -- stays launchable until gfx1250 is tuned
natively. Gluon configs are never copied across arches. Conv's section
carries the variant-aware walk from #4869 unchanged (its M-bounds renamed
CONV_STANDARD_M_BOUNDS to coexist with GEMM's).
Boss2002n added a commit that referenced this pull request Aug 27, 2026
utils/ reorganized for reuse: ONE config_utils.py absorbs core.py and the
gemm/conv/mhc/tuned config_utils modules, laid out in five clearly-banded
sections (shared core, GEMM, conv, MHC, tuned kernel entries) so each op
family stays visibly separate inside the single import surface. The
duplicate utils/mha_kernel_utils.py is dropped in favor of the byte-equal
utils/_triton copy; every importer tree-wide is updated. gmm_common stays
put under utils/ so tests and benchmarks keep importing a public module.

resolve_config_dir() gets the contract the #4947/#5019 reviews asked for,
hardened: an Args/Returns/Raises/Example docstring; every argument
validated against a whitelist with assertions (op, config_name and the
programmer-written arch= override must match strict identifier patterns,
the driver-derived running arch tolerates vendor formats but must stay
path-safe, backend is checked against ("triton", "gluon")) so a bad value
fails closed instead of resolving an escaped or wrong directory; a single
str return value (the unused name_prefix tuple element is gone, no caller
unpacks a tuple); and fully deterministic behavior -- the directory is
built, never probed, and the legacy escape hatches are deleted, so nothing
silently returns a bogus last candidate: a missing required table fails
loudly in load_config_json() naming the exact nested path, and optional
tables pass required=False.

The MOE dispatch loaders now declare their backend per dispatch path
(a8w4's triton path resolves backend="triton", the gluon paths "gluon";
a4w4 is gluon-only), and gfx1250 gains byte-identical triton seeds copied
from gfx950 (GEMM-AFP4WFP4, moe A8W4, HSTU fwd/bwd) so the one resolution
the caller-declared policy changes -- triton GEMM-AFP4WFP4 on gfx1250 no
longer borrowing gluon tables -- stays launchable until gfx1250 is tuned
natively. Gluon configs are never copied across arches. Conv's section
carries the variant-aware walk from #4869 unchanged (its M-bounds renamed
CONV_STANDARD_M_BOUNDS to coexist with GEMM's).
JiaoliangYu added a commit that referenced this pull request Aug 28, 2026
…error check (#5075)

* [FlyDSL] gfx942 a16wi4: pack f32->bf16 with lshr-16 instead of scalar (#5017)

* [FlyDSL] gfx942 a16wi4: pack f32->bf16 with lshr-16 instead of scalar truncf
v_cvt_pk_bf16_f32 is gfx950-only. After #4646 the gfx942 int4 fallback used
f32.to(bf16)/truncf, which is much more VALU than the old moe_gemm_2stage
shift-pack. Same nibble order; gfx950 packed convert and MXFP4 are unchanged.

* [FlyDSL] Clarify gfx942 a16wi4 upconvert comments

* ci: allow multigpu label to trigger tests (#5008)

* [HIP] [CK] [MoE] Added Gelu with tanh approx for CK XDL 2-stage MoE (#4620)

* [MoE] Added Gelu with tanh approx for CK XDL 2-stage MoE

* applied copilot's comment for str2ActivationType

* Dropping cross-activation CK configs

* Applied Ying comment

* Added block for run_1stage for unsuported activations

* fixed test

* [Triton] Move attention configs to nested layout and unify their resolution (#5019)

Relocate 14 attention config files from the flat arch-prefixed layout to
configs/<arch>/triton/attention/<d_type>/DEFAULT.json - mha, extend_attention,
mla_decode_rope, hstu_attn_fwd and hstu_attn_bwd - retiring configs/hstu_attn/.
The redundant -DEFAULT suffix is dropped from directory names (the file is
already DEFAULT.json), matching the chunk_delta_attn precedent. The six reader
modules resolve through the shared resolve_config_dir("attention", ...) probe
instead of hand-built paths. LEANATTN is not migrated: upstream removed the
lean_atten kernel and its config.

* [Triton] Migrate the GMM tuned configs to the nested layout (#5020)

Move configs/<arch>-GMM.json (gfx942, gfx950, gfx1250) to
configs/<arch>/triton/gmm/gmm/DEFAULT.json and point the reader at it.
GMM gets its own op directory instead of folding under gemm/. The
doubled gmm/gmm is just the <op>/<d_type> layout: the op is "gmm" and
the family's config name is "GMM", so _dtype_dir() yields "gmm" too.

_triton_kernels/gmm.py now resolves the directory through the shared
resolve_config_dir("gmm", "GMM", backend="triton") probe and loads
DEFAULT.json from it; arch_info and AITER_TRITON_CONFIGS_PATH are dead
there and are dropped. No legacy_dir is passed - the files move and the
loader flips in this one commit, so every revision resolves.

* [Triton] Move MOE tuned configs to the nested layout (#5022)

Move the three remaining MOE tuned configs from the flat configs/moe/
directory into configs/<arch>/<backend>/<op>/<d_type>/:

  moe/gfx950-A8W4.json  -> gfx950/triton/moe/a8w4/DEFAULT.json
  moe/gfx1250-A8W4.json -> gfx1250/gluon/moe/a8w4/DEFAULT.json
  moe/gfx1250-A4W4.json -> gfx1250/gluon/moe/a4w4/DEFAULT.json

The backend directory follows the dispatch path the table actually feeds,
not the arch: gfx950's a8w4 table is keyed bm<block_m>_n<N>_k<K> and is
read by the Triton path, while both gfx1250 tables are bucket-keyed and
read by the Gluon path. So the a8w4 family spans backends and a4w4 is
Gluon-only.

These three are all that is left of configs/moe/: PR #4833 removed the
rest of the legacy MOE stack (utils/moe_config_utils.py, the fused
sigmoid-top1 routing kernel, the moe_op/moe_op_e2e/mxfp4 variants and
every configs/moe/*-MOE-*.json), so this completes the directory.

The two surviving loaders are rewired onto the shared probe in the same
commit. _get_a8w4_dispatch() and _get_a4w4_dispatch() now resolve their
directory with resolve_config_dir("moe", "<A8W4|A4W4>") and read
DEFAULT.json from it, instead of hand-building an arch-prefixed path
under configs/moe/. Neither call passes backend=: because the backend
differs per arch for the same family, pinning one would make the other
arch's file unreachable. The documented probe order -- nested triton,
then nested gluon -- picks whichever directory the running arch ships.

a4w4 also moves off its private os.path.exists + json.load pair onto
load_config_json(..., required=False), matching a8w4; both still return
{} when no tuned file is shipped for the arch, so the safe-default
fallback paths are unchanged.

resolve_config_dir() lives in utils/gemm_config_utils.py and is added by
the config-unification branch -- merge that one first.

* [Doc][Skill] port udpated flydsl kernel code cleanup skill (#5051)

* [Triton/Gluon] MOE a8w4 cudagraph updates (#5037)

* [Triton] Remove legacy flat-layout support from config resolution (#4948)

* [Triton/Gluon] Move gluon gemm_a8w8 kernel into _gluon_kernels/gfx950 (#4866)

* [Triton] Migrate conv configs to the nested arch/backend layout (#5018)

Move all 59 flat configs/conv/<arch>-<CONFIG_NAME>.json files to
configs/<arch>/triton/conv/<d_type>/DEFAULT.json, the layout GEMM already
uses, and point _conv_config_path() at the shared resolve_config_dir()
probe instead of building the legacy path by hand. This picks up the ten
tables #4869 added (CONV-PREPACK on all seven arches, CONV-3X3-NCHW on
gfx1100/gfx1151/gfx1201) alongside the original 49.

The renames and the loader flip land in one commit so every revision
resolves conv configs from exactly one layout: no legacy_dir fallback is
needed and bisect stays clean. File contents are untouched (pure renames).
_conv_config_path() is the single choke point, so get_conv_config(),
has_conv_config(), conv_config_uses_exact_routes() and
has_exact_conv_config() all pick up the nested path; the variant-aware
four-tier walk, STANDARD_M_BOUNDS and the lru_caches are untouched.

* Tune MoE GEMM A8W8 blockscale (#5028)

* [Triton] Migrate MHC configs to the nested arch/backend layout (#5021)

Move all 15 flat configs/<arch>-MHC_*.json files to
configs/<arch>/triton/mhc/<d_type>/, keeping the C=<n> specialized file
stems and naming each family default DEFAULT.json, and rewire
mhc_config_utils onto the shared resolve_config_dir() probe. The
documented gfx942 fallback retry resolves through the probe's arch=
override (added by the legacy-removal PR, which merges first); the
C-bucket walk and _FALLBACK_DEV semantics are unchanged.

The renames and the loader flip land in one commit so every revision
resolves MHC configs from exactly one layout, and this branch touches
no shared resolver code.

* [Gluon] add bench for mxfp8 GEMM (#5029)

* [FlyDSL] One-stage split-K for the a8w8 preshuffle GEMM (#5007)

* [FlyDSL] one-stage split-K for the a8w8 preshuffle GEMM

Fold the split-K reduction into the GEMM launch: every split publishes an
fp32 partial, and the last one to arrive at the tile's semaphore reduces
and converts in the same kernel, so split-K costs one launch rather than
two.

The partials cross CTAs that may sit on different XCDs, each with its own
L2, so they have to reach a common point. Doing that with an agent-scope
fence costs a whole-L2 buffer_wbl2 per CTA plus a buffer_inv on the
reader, which also evicts the A/B tiles every other in-flight CTA is
still reading -- measured 2-3x slower than not splitting at all. Marking
just these accesses sc0|sc1 writes them through and leaves L2 alone.
That turns split-K from a 18-180% regression into a 9-50% win over
k_split=1.

k_split == 1 is untouched: same kernel, same cached stores, and its
output is bitwise identical to main across the shapes checked.

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

* [FlyDSL] drop the k_split == 2 special case

It had its own epilogue: whichever split arrived first published its fp32
fragment, and the second spun on a ready flag, kept its own fragment in
VGPRs and wrote the final tile -- saving one workspace plane and one
round trip.

It does not pay for itself. Spinning is slower than just going through
the generic path: 4.7 vs 5.8 us at 1x576, 18-27% across the six shapes
measured. Removing it also drops two fragments, two copy atoms, the
doubled semaphore, and the split-plane special cases in the launcher and
the AOT pre-compile.

k_split == 2 now takes the same path as every other split count, which
also fixes the per-split workspace offset: it was guarded on
split_k > 2, so a k_split == 2 launch routed through the generic path
would have had every split write the same plane.

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

* [FlyDSL] fix the split-K keyword the tuner passes to the launcher

The tuner called the preshuffle launcher with k_split=, but that launcher
names the argument split_k= (matching the hgemm split-K path it sits next
to). Every flydsl preshuffle candidate therefore raised TypeError.

The tuner records a raising candidate as rejected rather than as an error,
so the run completed, kept only the 8wave candidates, and picked ck or
cktile for four shapes that flydsl had previously won -- a result
indistinguishable from a legitimate tuning outcome.

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

* [FlyDSL] trim the split-K diff to what split-K needs

Three things that were not split-K:

out_dtype grew an fp32 branch and a raise. Nothing needs it -- fp32 is the
type of the *partial*, which "Float32 if split_k > 1" already covers, and
the final output is still bf16/fp16. The bias element-type change existed
only to feed that branch. Both are back to main's two-case form.

The K-tile index has to gain a bid_z offset, which is genuine, but the
name k_tile_base pushed several one-line fx.copy calls past the line
limit and a trailing comma pinned others open, so a one-token change read
as +5 -1. Renaming to k_off and dropping the magic trailing commas keeps
them one-liners.

The copy atom for the output no longer branches on out_elem_bytes; it
picks the op from split_k directly.

Kernel diff: +201 -22 -> +168 -21, with no behaviour change. k_split == 1
still compiles to bytes identical to main.

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

* [Config] retune Kimi-K3 a8w8 M<=32 with one-stage split-K

29 rows, all M <= 32: 21 stay on flydsl with a better config and 8 move
from ck to flydsl. 20 of them use split-K, mostly k_split=7.

The tuner proposed 37 rows. Each changed row was then re-measured old
config against new on an idle GPU, and the 8 that were actually slower
there were kept at main's value -- the tuner picks its winner while four
GPUs are saturated, and for shapes where several configs sit within noise
of each other that choice does not survive on an idle card. Nearly all of
them were k_split=2 at N=6400, which lost 7-13%.

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

* [FlyDSL] hoist the partial store out of the split_k branch

Both arms opened with the same fx.copy; only what follows it differs.

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

* [FlyDSL] address review: semaphore dtype, buffer lifetime, reduce guard

The k_split == 1 path passed an empty bf16/fp16 tensor in the semaphore
slot while the AOT pre-compile passed an empty int32 one. dtype is part
of FlyDSL's executable cache signature, so every non-split-K preshuffle
kernel missed its AOT entry and JIT-compiled at first call -- a
regression across all existing tuned configs, not just split-K. Both
sides now pass int32.

The split-K buffers were cached per (m, n, tile, k_split). m is in the
key, so a server sweeping batch sizes grows the cache without bound, and
an eviction frees memory whose address a captured CUDA graph still holds.
k_split_candidates only proposes split-K while the tile grid is under one
CTA per CU and caps k_split * tile_count at four per CU, which bounds
tile_count below CU_NUM and the workspace at 4 * CU_NUM * tile_m *
tile_n floats -- so the buffers are now fixed-size and keyed on
(device, stream) only, the way _get_split_k_tensors already does it, with
a capacity check for anything that would exceed the bound.

The reduce derives its vector count as tile_n // 4 and would have dropped
the tail columns for a tile_n that is not a multiple of 4; the comment
claimed the invariant but nothing enforced it. Now rejected at compile
time.

The semaphore reset was a plain cached store while every other cross-CTA
access in that block carries sc0|sc1. It is the same cross-XCD handoff,
one launch later, so it writes through too.

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

* [FlyDSL] guard M against the layout bound the kernel assumes

The kernel views A and C through layouts with a hardcoded 65536 rows, in
three places, with nothing on the host stopping a larger M from indexing
past them. Named the bound, used it at all three sites, and rejected an
out-of-range M in the launcher with a message that says why.

gemm_kernels keeps its own copy of the literal because that module has to
import without FlyDSL present; a test asserts the two agree.

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

* [FlyDSL] cut the comments back to what the code cannot say

57 added comment lines down to 26. Dropped the ones restating the line
below them -- what k_off is, which path split_k > 1 takes, that the
partial store publishes a partial -- and shortened the rest.

What is left is the reasoning that is not recoverable from the code: why
the partials cannot use an agent-scope fence, why the buffers are fixed
size rather than shape-keyed, why the semaphore dtype has to match the
AOT side, why _REDUCE_VEC is 4, and why the k_split candidates are
enumerated per shape.

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

* Extract the flydsl split-K reduce into a reusable copy-atom epilogue

Move the one-stage split-K reduction out of preshuffle_gemm into
splitk_epilogue.splitk_reduce_epilogue, with the output element class as its
only dtype knob so other GEMMs can reuse it.

The reduce now goes through copy atoms and a buffer-tensor descriptor instead
of raw buffer_ops: make_layout_tv gives each thread 4 contiguous columns, so
the loads stay dwordx4 and the stores dwordx2, and the descriptor bounds cover
the ragged-M tail. Resetting the semaphore with atomic_add(-split_k) rather
than a plain store also drops a next-launch increment race.

Verified on gfx950: rel_err matches k_split=1 for M in {1,8,64} x k_split in
{1,2,7,14}, CUDA-graph replays clean, and the k_split=1 output hashes are
identical to origin/main. Over the 20 tuned Kimi-K3 split-K shapes the reduce
is 0-3% faster than the buffer_ops version.

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

* Let --splitK gate the FlyDSL candidates as it does the other backends

The ck, cktile and asm task builders all take useSplitK and collapse the
split-K dimension to a single splitK=0 candidate when it is off. The FlyDSL
builder never received the flag, so it generated k_split candidates
unconditionally: on the Kimi-K3 shape set that is 98488 extra candidates on
top of 99136, roughly double overall and 2.5-3.1x over M in 1..128.

It also made the flag useless as a switch. Split-K wins often enough at
small M that 20 of the 176 tuned rows are split-K winners, so a run without
--splitK still produced split-K configs.

---------

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

* [FlyDSL] Retune GLM5.2 mxfp4 MoE and fix a scale-view cache leak (#5045)

* [FlyDSL] Retune GLM5.2 mxfp4 MoE and fix a scale-view cache leak

Retune all 64 GLM5.2 shapes (model_dim=6144, inter_dim 256..2048, E=257,
topk=9) for gfx950. 27 shapes move to the coupled flydsl_mxmoe port, which
the previous config only reached on 5 rows.

Measured through the production fused_moe path, each shape timed on one GPU
under both configs: median +6.7%, mean +8.6%; 44/64 faster by >1%, 6 slower
(worst -3.5%). Small batches gain most (token<=64 median +11.4%).

Six shapes (2048/256, 4/1024, and 2/16/64/128 at 2048) are left on their
existing main entries rather than retuned.

Two fixes fell out of the tuning runs:

_mxfp4_scale_u8 was wrapped in lru_cache(maxsize=2048). Its body is a bare
.view(), so the memo buys nothing, but tensors hash by identity: every
per-call intermediate scale misses and is then pinned by the cache. A tuning
sweep leaked ~0.75 GiB per timed iteration and exhausted a 288 GiB card.

v2_stage1_dequant_cosine_err looped per sorted row, costing one .item()
sync each -- ~295k syncs per timed candidate at token=32768/topk=9. Now
batched in chunks, which bounds the int64 gather in mxfp4_to_f32 while
keeping the equal-weight average over rows.

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

* [FlyDSL] Default FMoE tuning to FlyDSL v2 and update GLM5 FP4 layout configs

* fix black test

* Emit the non-f4out AOT job for _f4out mxmoe stage-2 rows

An `_f4out` GEMM2 kernel only really runs the mxfp4-out path when both
gates are open: AITER_MXFP4_INTERMEDIATE, and the shape check in
fused_moe (`D_HIDDEN == 7168`). Otherwise `_f4out` is stripped from the
kernel name and the plain kernel launches instead.

The AOT generator skipped such rows outright, so it never pre-compiled
the kernel that actually launches. GLM5 is D_HIDDEN=6144, so every
`_f4out` row there falls back -- and the config only survived because an
unrelated row happened to name the plain kernel and seed the same cache
entry. Retuning that row to `_f4out` removed the last such seed and CI
hit `FLYDSL_RUNTIME_RUN_ONLY=1 but no usable AOT cache for launch_gemm2`
on token=16384, inter_dim=512.

Emit the fallback job unconditionally, plus the f4out one when
AITER_MXFP4_INTERMEDIATE is set.

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

---------

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
Co-authored-by: charlieguo1106 <cguo@amd.com>

* [UT] Support a4w4 in test_mega_moe (#5052)

* support a4w4 in test_mega_moe_gfx1250

* support 64K

* [Triton/Gluon] [ASM] [HIP] add mla v4 prefill asm kernel (#4926)

* Add MLA v4 sparse prefill asm support

Integrate the gfx1250 MLA implementation and consolidate sparse prefill correctness and performance tests.

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

* Update op_tests/test_pa_sparse_prefill.py

Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>

* Apply black formatting to test_pa_sparse_prefill

Pure reformat, no behaviour change. Fixes the failing black CI job
(black[colorama]==26.5.1, default line length).

* Fix asm candidate reference in test_pa_sparse_prefill

The asm candidate passed split["ref"] -- the raw input dict -- where
checkAllclose expects the reference tensor, so the first asm comparison
died with:

  TypeError: isclose(): argument 'other' (position 2) must be Tensor, not dict

meaning the asm path could never run. Compute the fp8 reference the same
way the opus fp8 candidate above it does.

* Report per-row nnz and default the CLI to the asm comparison sweep

Two test-driver changes:

* nnz_prefix/nnz_extend columns now report per-row nnz instead of the
  pool-wide total, so they match the --nnz-prefix/--nnz-extend asked for
  rather than scaling with N. total_nnz still carries the full count --
  the TFLOPS/TB-s figures need the real work done.

* CLI defaults now describe the three-way opus/triton/asm comparison out
  of the box: N in [512, 1024, 2048, 4096] x nnz_prefix in
  [256, 1024, 4096, 8192, 16384] x nnz_extend 128, at H_Q=128 fp8 (the
  only shape the asm candidate registers for). --mode/--total_pages
  default empty so the unrelated mode sweep stays off unless asked for.
  Every flag still overrides. Pytest coverage is unaffected: it reads
  _PYTEST_SHAPES/_PYTEST_MODES, not argparse.

* Accept an over-allocated CSR indptr in mla_sparse_prefill

check_csr required indptr->numel() == T+1 exactly. Decode reuses this
kernel with the extend region empty and sizes its CSR row-pointer buffers
once at [max_batch+1], launching with the live batch, so numel > T+1 is
the normal case there rather than a mistake -- and the exact test rejected
it outright.

The kernel reads indptr[0..T] and nothing past it, so the extra tail is
inert: verified bit-identical output against the exactly-sized call. An
undersized indptr is still rejected.

Trade-off: an indptr built for a different T is no longer caught here.
Separating that from the legitimate case needs device data (indptr[T]
against the indices length), i.e. a sync per call. Callers that can slice
to [:T+1] should.

* Fix int32 overflow in sparse prefill query offset

`_sparse_attn_prefill_kernel` derived `query_idx` from `tl.program_id(0)`,
which Triton types as int32. The q/out addresses are computed as
`query_idx * q_stride_t` and `query_idx * out_stride_t`, and in the V4
layout that stride is `num_heads * head_dim` = 128 * 512 = 65536. The
product therefore leaves the int32 positive range at `query_idx >= 32768`
and wraps to a negative offset, so the kernel reads and writes outside the
q/out allocations.

Observed as NaNs followed by a hard GPU page fault:

    Memory access fault by GPU node-2 ... Reason: Page not present

Verified on gfx1250 with a fixed-pattern sparse prefill case
(H=128, D=512, pool=16384, nnz_prefix=256, nnz_extend=128):

    N=32768  before: clean   (largest size that still fits int32)
    N=32769  before: fault   after: nan=0 inf=0
    N=65536  before: fault   after: nan=0 inf=0

Promoting `query_idx` to int64 moves both offsets to 64-bit address
arithmetic. This mirrors the existing `slot_off` cast a few lines below,
which already handles the same class of overflow on the pool index; the
difference is that the wrapped pool offset stays inside the allocation and
reads silently, while this one faults.

---------

Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>

* [HIP] FIX MLA the nhead fold error for cp round robin (#4964)

* fix the nhead fold error for cp round robin

* fix the split test

* support varlen

* gqa96 qseqlen<6 not fold

* [HIP] [Bugfix] Fix DSV4 FP4 KV-cache scattered row writes (#5034)

* fix(dsv4): scatter FP4 KV cache by row-local offset

Signed-off-by: AMD-yanfeiwang <yanfei.wang@amd.com>

* test(dsv4): remove specialized KV-cache regression

Keep the bug fix focused without carrying a narrow special-case test.

---------

Signed-off-by: AMD-yanfeiwang <yanfei.wang@amd.com>

* [ASM] [HIP] 1x32 mxfp4 asm kernel (#4890)

* 1x32 mxfp4 asm kernel

* Update tuned config

* Upate 1x32 kernel to embedd X quant

* Drop the standalone MXFP4 X quant pre-pass plumbing

The FLAT MXFP4 kernels dynamic-quantize X in-kernel, so the host-side
pre-pass entrypoint, its Python binding and the test helper have no
caller left. Removing them also restores the per_1x32 scale-sorting
condition, which still tested a pre-pass flag that no longer exists.

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

* Tune 1x32 kernel

* Fix 1x32 race condition for O buffer clearning

* SImplify zero protocall and bind it to TG0 always

* Fix lm_eval utter failure with 1x32 kernel

---------

Co-authored-by: Sergey Solo <ssolovye@amd.com>
Co-authored-by: Cursor <cursoragent@cursor.com>

* [HIP] update MHA CPP reademe (#4874)

* update the supported arguments configuration

* update the perf data

* update the image

* fix the log

* fix

* benchmark_fwd support opus kernel

* add opus perf data

* perf data

* fix the comment

* fix

* Test FFM bringup on MI250 build runner (#5071)

* perf(gfx1250): drop a16w16's 4 GiB pre-check, fail on wrong results

The bench skipped any a16w16 shape whose largest operand passed 4 GiB, on
the stated grounds that "the heuristic refuses" it. That reads the guard
too broadly.

opus_dispatch_a16w16_gfx1250 (opus_gemm_arch_gfx1250.cuh:150-183) searches
the tuned table first and returns on a hit. check_shape_4g runs only after
that misses, on the way to the split-K heuristic kid, and it is that kid's
launcher that builds the 32-bit gmem descriptors. A tuned 4wave_wl_co
winner never reaches the check: per gen_instances_gfx1250.py:770-778 the
pipeline "builds no gmem descriptor at all" and clamps every dimension
through TDM descriptors instead. So 4 GiB bounds one fallback path, not
a16w16, and a pre-check in Python keeps skipping shapes that tuning has
already made runnable. Removed; the kernel raises if it must, and the
exception is recorded as a row.

The 20260827 sweep shows what the fallback costs. At N=129280 the tuned
4wave_wl_co kid does M=512 in 449us (2112 TFLOPS); M=1024 has no tuned
winner, drops to split-K, and takes 3343us (568 TFLOPS) -- 7.4x slower for
2x the work. 11 of 60 shapes hit a 4wave_wl_co kid; the rest are split-K,
so most of the low numbers in this table measure tuning coverage rather
than the hardware. Widening that coverage is a job for
csrc/gemm_a16w16/gemm_a16w16_tune.py --libtype opus, not for this file.

Worse, split-K is not just slow at the top of the range: all four M=65536
shapes came back err=0.96-0.99, an unrelated result, while every other row
was 0 or ~1e-5. None of them trip the 4 GiB guard (M*K*2 = 896 MB, M*N*2
<= 256 MB), and the UT neither raises nor warns -- it returns the ratio and
prints a number. The sweep reported them as data. a16w16 now checks the
returned ratio against _A16W16_MAX_ERR and calls _note_failure, so a silent
miscompare shows up in the failed-op list.

a16w16 also gets its own M list. The global sweep jumps 2048 -> 65536, so
the prefill chunk sizes were never measured on the BF16 linears; _A16W16_MS
adds 4096/8192/16384 and AITER_BENCH_TOKENS still overrides it.

The lm_head cap stays. It is a statement about what DSv4 runs -- one row
per sequence -- not about what the kernel can do, and its comment no longer
leans on the 4 GiB number.

Separately, put a8w8_blockscale back in --dsv4 and correct its note. The
note blamed #4773's gluon tuning rows for the make_llir crash. The real
cause is the UT's extra "ck strided x_scale" check
(test_gemm_a8w8_blockscale.py:120), added by #4406 and gated on
ck_preshuffle alone. The mxfp8_128 path declares its layout with
is_x_scale_transposed=True and never reads the stride, so a strided x_scale
tests nothing there and only gives triton a specialization that fails to
compile. A/B with that line as the only variable, over a 162-case matrix
(27 default M x six (n,k)): case 2 before it dies, case 160 after -- M=16
and M=64 included, which is what #4773 covers. Fixing it properly is
upstream's call; meanwhile _A8W8_BLOCKSCALE_TOKENS starts at 1024, clear of
the M that reach those rows. Verified on gfx1250-atom--20260827-ubench:
36/36 cases, err=0, 2207-7003 TFLOPS.

---------

Signed-off-by: AMD-yanfeiwang <yanfei.wang@amd.com>
Co-authored-by: msaffari-amd <msaffari@amd.com>
Co-authored-by: Xin Huang <Xin.Huang@amd.com>
Co-authored-by: Alexandra Sidorova <asidorov@amd.com>
Co-authored-by: Satya Nikhil Kodukula <nikhil.kodukula@gmail.com>
Co-authored-by: Felix Li <felix.li@amd.com>
Co-authored-by: Lukasz Burzawa <lukasz.burzawa@amd.com>
Co-authored-by: Vinayak Gokhale <vinayak.gokhale@amd.com>
Co-authored-by: Nidal Danial <81209936+nidal567@users.noreply.github.com>
Co-authored-by: Shao-Chun Lee <Shao-Chun.Lee@amd.com>
Co-authored-by: XiaobingZhang <xiaobingzhangupc@gmail.com>
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
Co-authored-by: charlieguo1106 <cguo@amd.com>
Co-authored-by: yanboshao <yashao@amd.com>
Co-authored-by: junxiaguo <JunXia.Guo@amd.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>
Co-authored-by: minmengdie <memin@amd.com>
Co-authored-by: AMD-yanfeiwang <yanfei.wang@amd.com>
Co-authored-by: Sergey Solovyev <sergey.solovyev@amd.com>
Co-authored-by: Sergey Solo <ssolovye@amd.com>
Co-authored-by: Yu <jiaolyu@amd.com>
brunomazzottiamd pushed a commit that referenced this pull request Sep 8, 2026
…DNA3) and gfx1150 (RDNA3.5) (#5246)

* [Triton] [Conv2D] Add Conv2D configuration files for gfx1101 and gfx1150

Nine tables per architecture in the nested layout, tuned on hardware. Both
architectures currently have none, so every Conv2D family is a hard
FileNotFoundError from load_config_json(required=True) rather than a fallback.

gfx1101 (RDNA3, RX 7800 XT): 190 entries, 20 tuned on the card, 170 inherited
byte-identical from gfx1100.
gfx1150 (RDNA3.5, Radeon 890M): 186 entries, 100 tuned, 86 inherited from
gfx1151.

Entry counts cover every section - shapes, variant pins, M_LEQ_* buckets and
any - not just shapes.

The structure of every file matches its donor exactly: same sections, same key
set per section, same route_exact_only. Routing behaviour is therefore unchanged
and only the values differ. In particular CONV-3X3-NCHW ships only the shape
keys the donor already has, because under route_exact_only the key set is what
_nchw_direct_is_profitable uses to choose between DIRECT_NCHW_3X3 and
CBLOCKED_NCHW, and that comparison was not part of these measurements.

Every candidate was checked for correctness before being timed - rel_l2 < 1e-2
against torch.nn.functional.conv2d, the bar op_tests/triton_tests/conv uses -
and each family was driven through its own public entry point rather than
through the router.

Follows the invitation in
#4869 (comment)

Signed-off-by: Martin Domanský <ragua@email.cz>

* [Triton] [Conv2D] Add gfx1101 and gfx1150 to the conv test SUPPORTED_ARCHS

Without this the two new tables are never exercised by op_tests/triton_tests/conv.

Note that the repository's CI has no RDNA machine, so this does not by itself
cause the tables to be tested anywhere - it only makes the tests runnable on
hardware that has them.

Signed-off-by: Martin Domanský <ragua@email.cz>

* [Triton] Add the two new arches to the config layout table

configs/CLAUDE.md documents which architectures have a config tree. #5085
rewrote that document for the nested layout, so it is maintained and a new arch
belongs in it.

Separate commit from the data on purpose: it is your document, and a doc change
should be reviewable on its own rather than folded into eighteen JSON files.

Signed-off-by: Martin Domanský <ragua@email.cz>

* [Triton] [Conv2D] Drop CONV-3X3-NCHW from this PR

That table is the optional ninth family and adding it is a routing change on
these two architectures, which the rest of this PR is not.

_nchw_direct_is_profitable returns False when has_conv_config("CONV-3X3-NCHW")
is False, so today every 3x3 shape on gfx1101 and gfx1150 takes CBLOCKED_NCHW.
Shipping the table with route_exact_only moves 7 shapes on gfx1101 and 4 on
gfx1150 onto DIRECT_NCHW_3X3, and none of those 11 has a direct-against-cblocked
measurement on either card.

The tables had already been trimmed to the donor's key set so they would carry no
routing change relative to gfx1100 and gfx1151. That was the wrong baseline: the
comparison that matters is against what these cards do now, which is no table at
all.

gfx1200, gfx1250, gfx942 and gfx950 ship no CONV-3X3-NCHW either, so the
remaining eight families are the mandatory set and this PR now carries values
only.

Counts after this commit:

  gfx1101   182 entries, 16 tuned on the card, 166 inherited from gfx1100
  gfx1150   181 entries, 97 tuned on the card,  84 inherited from gfx1151

The table returns in a follow-up that measures both routes per key.

Signed-off-by: Martin Domanský <ragua@email.cz>

---------

Signed-off-by: Martin Domanský <ragua@email.cz>
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.

4 participants