[Triton/Gluon] Moe routing optimizations

[lburzawa]

Motivation

Optimize stores in routing kernels to reduce number of pids writing to the same address and prevent cash thrashing. Issue originally pointed out by @amd-hhashemi.

Technical Details

Issue only the stores that are necessary. Pass data between expert metadata kernels using tl.gather instead of store->load sequence.

Test Plan

UTs and performance.

Test Result

UTs pass. Significant performance improvement on cache coherent driver on MI450.

Submission Checklist

• Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

[github-actions]

🏷️ 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 5038 --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.

[Copilot]

🟡 Changes recommended

Address zero-row token-offset initialization and provide reproducible performance benchmarks.

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

Pull request overview

Optimizes MoE routing metadata propagation and reduces redundant stores between routing stages.

Changes:

• Uses tl.gather to pass tile offsets directly.
• Masks histogram loads and limits metadata writes.
• Updates fused and non-fused routing paths.

File summaries

File	Summary
aiter/ops/triton/_triton_kernels/moe/moe_routing/routing.py	Consumes returned tile offsets in routing kernels.
aiter/ops/triton/_triton_kernels/moe/moe_routing/expt_data.py	Refactors metadata computation; zero-row inputs may leave token offsets uninitialized, and benchmark evidence is missing.

Review details

Suppressed comments (1)

aiter/ops/triton/_triton_kernels/moe/moe_routing/expt_data.py:35

• Calling tl.gather directly here makes this routing kernel depend on a Triton build that exposes that API. The repository already detects hasattr(triton.language, "gather") and provides a fallback in aiter/ops/triton/_triton_kernels/chunk_delta_attn/intra_attn.py; on a system-Triton build without it, the first routing compilation fails instead of falling back. Author must route this through a guarded helper/fallback or explicitly enforce the minimum Triton version.

        tile_start = tl.gather(tile_starts, expt_id, 0)

• Files reviewed: 2/2 changed files
• Comments generated: 2
• Review effort level: Lite

---

💡 Configure MCP servers for context-aware, tailored reviews. Learn more in the docs.