Skip to content

[Triton/Gluon] [JIT] # RDNA3 fp8 enablement + tuned DSV4 A8W8 block-scale GEMM configs - #5217

Merged
Boss2002n merged 4 commits into
ROCm:mainfrom
amd-xavierwang:gfx1100-blockscale-tuning-wip
Sep 17, 2026
Merged

Boss2002n merged 4 commits into
ROCm:mainfrom
amd-xavierwang:gfx1100-blockscale-tuning-wip

Conversation

@amd-xavierwang

@amd-xavierwang amd-xavierwang commented Sep 2, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

  • gfx11 has no default configs nor shape specific config files for fp8-gemm-blockscale in aiter now.
  • gfx11 uses fnuz for dtype which should be changed to torch.float8_e4m3fn.
  • the GRID_MN is a tl.constexpr which makes Triton to trigger re-compile for its dependencies, M and N. For LLM, cold compiling for each M in GEMM is not accepted.

Changes

1. aiter/ops/triton/utils/types.py — gfx11 fp8 dtype

Add gfx1100, gfx1101, gfx1102, gfx1151 to the non-fnuz arch tuple in both
get_fp8_dtypes() and get_fp8_e4m3_dtype() → e4m3 resolves to
torch.float8_e4m3fn on gfx11.

2. aiter/ops/triton/configs/.../gemm_a8w8_blockscale/ — tuned configs

Six JSON files in the nested layout (#4927):

file shape (N, K) role
DEFAULT.json — fallback covering all (N, K)
GEMM-A8W8_BLOCKSCALE-N=1536-K=4096.json 1536, 4096 DSV4 TP=8 linear
GEMM-A8W8_BLOCKSCALE-N=512-K=4096.json 512, 4096 DSV4 TP=8 linear
GEMM-A8W8_BLOCKSCALE-N=4096-K=1024.json 4096, 1024 DSV4 TP=8 linear
GEMM-A8W8_BLOCKSCALE-N=4096-K=256.json 4096, 256 DSV4 TP=8 linear
GEMM-A8W8_BLOCKSCALE-N=8192-K=1024.json 8192, 1024 DSV4 TP=8 linear

3. aiter/ops/triton/_triton_kernels/gemm/basic/gemm_a8w8_blockscale.py — remove M-dependent GRID_MN

GRID_MN causes M, N to become constexpr, causing a new JIT compile on every new M. This is not acceptable for LLM inference. The only usecase, remap_xcd(pid, GRID_MN) is also dead code.

4. RDNA3/3.5 fp8 enablement

file change
aiter/ops/triton/utils/types.py add gfx1150 to the non-fnuz arch tuple → e4m3 resolves to torch.float8_e4m3fn
aiter/utility/dtypes.py add "gfx1150": {"fp8": torch.float8_e4m3fn} to defaultDtypes (see #4764 )
aiter/jit/utils/chip_info.py add gfx1150 to _LDS_CAPACITY_BYTES at 64 KiB/WGP (also fills the missing gfx1101/gfx1102/gfx1200 entries)

Exercised through vLLM DeepSeek-V4-Flash serving on gfx1100 (TP=8, fp8 KV cache). Inference runs after adapting Aiter's backend.

Test Result

Metric Baseline aiter blockscale Δ
Output throughput (tok/s) 21.09 30.96 +46.8%
Benchmark duration (s) 1006 660.7 −34.3%
Median TPOT (ms) 200 113 −43.5%
Median ITL (ms) 421 191 −54.6%
# tuned config: GEMM-A8W8_BLOCKSCALE-N=8192-K=1024.json
pytest op_tests/triton_tests/gemm/basic/test_gemm_a8w8_blockscale.py \
  -k "False and triton and 8192-1024"
# -> 6 passed, 6 skipped (gluon)

# tuned config: GEMM-A8W8_BLOCKSCALE-N=1536-K=4096.json
pytest op_tests/triton_tests/gemm/basic/test_gemm_a8w8_blockscale.py \
  -k "False-triton and 1536-4096"
# -> 6 passed

# DEFAULT.json fallback (no per-shape file for 4096x4096)
pytest op_tests/triton_tests/gemm/basic/test_gemm_a8w8_blockscale.py \
  -k "False-triton and 4096-4096"
# -> 7 passed

Submission Checklist

*** Honouring contribution to @saeid-rostami for identifying the default dtype and compilation regression.

@amd-xavierwang
amd-xavierwang requested a review from a team September 2, 2026 17:07
@github-actions github-actions Bot changed the title # [Triton/Gluon] fp8 e4m3 RDNA3 dtype fix + tuned DSV4 A8W8 block-scale GEMM configs [Triton/Gluon] # fp8 e4m3 RDNA3 dtype fix + tuned DSV4 A8W8 block-scale GEMM configs Sep 2, 2026
@github-actions

github-actions Bot commented Sep 2, 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 5217 --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.

This comment was marked as outdated.

@amd-xavierwang

amd-xavierwang commented Sep 2, 2026 •

Copy link
Copy Markdown
Contributor Author

Agreed. In addition, kpack and matrix_instr_nonkdim are essentially no-ops on RDNA. I will do the clean up with existing files. There are gfx1201 files in upstream with kpack as well. I can clean those up as well upon demand.

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

Three specialized configurations are never selected by the existing unit-test matrix.

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

Review details
  • Files reviewed: 7/7 changed files
  • Comments generated: 3
  • Review effort level: Balanced

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

Architecture-specific dtype behavior and GPU tuning require human confirmation, and one duplicated dispatch list remains.

Review details
  • Files reviewed: 7/7 changed files
  • Comments generated: 1
  • Review effort level: Balanced

Comment thread aiter/ops/triton/utils/types.py Outdated
get_fp8_dtypes / get_fp8_e4m3_dtype fell through to the fnuz fp8 types on gfx1100/1101/1102/1151, so fp8 e4m3 tensors from vLLM were mis-typed. Add the gfx11 archs to both, mirroring the defaultDtypes map (ROCm#4764).

Signed-off-by: xawang <xawang@amd.com>
Tuned DeepSeek-V4 TP=8 block-scale GEMM shapes for gfx1100 (RDNA3): N=1536-K=4096, N=512-K=4096, N=4096-K=1024, N=4096-K=256, N=8192-K=1024, plus DEFAULT fallback. Placed in the nested config layout (ROCm#4927).

Signed-off-by: xawang <xawang@amd.com>
Squashes the post-tuning cleanup for the gfx1100 A8W8 block-scale GEMM
and extends fp8 support to gfx1150 (RDNA3.5).

Remove M-dependent GRID_MN constexpr from the block-scale kernels.
GRID_MN was promoted to a tl.constexpr via @triton.heuristics, so its
value (a function of M) entered the Triton compile-cache key. Every new
M therefore cold-compiled a fresh kernel, even though GRID_MN's only
consumer was remap_xcd(pid, GRID_MN) whose return value was discarded --
a no-op on every arch (and inert on single-XCD RDNA anyway). Drop the
GRID_MN heuristic, constexpr param, repr entry, and the dead remap_xcd
call (plus its now-unused import) from both _gemm_a8w8_blockscale_kernel
and the preshuffle variant. M now reaches the key only through the
16-divisibility specialization, matching the one-binary-per-shape
behavior of the vLLM-native kernels. num_pid_m/num_pid_n are still
computed at runtime and the masked store handles arbitrary M/N, so this
is behavior-preserving. Verified cos=1.0 / relL2<=1e-4 across M in
{1,2,3,16,17,31,64,128,129,256,384,512}; with a fixed config it compiles
<=2 binaries for all 12 M values (was ~1 per distinct M).

DRY get_fp8_e4m3_dtype() to delegate to get_fp8_dtypes() rather than
duplicating the gfx arch list (per PR review). The returned e4m3 dtype is
identical, so no behavior change.

Drop the deprecated kpack field from the tuned gfx1100 block-scale
configs; it is ignored by the current kernel and its presence tripped
config-schema validation.

Add gfx1150 (RDNA3.5) to fp8 dtype selection (OCP e4m3fn, like the other
RDNA archs), to the default-dtype table, and to the LDS-capacity table
(64 KiB per WGP) so the arch is dispatchable end-to-end.

Signed-off-by: xawang <xawang@amd.com>
@amd-xavierwang
amd-xavierwang force-pushed the gfx1100-blockscale-tuning-wip branch from a5e8eb9 to ee849d6 Compare September 10, 2026 16:57
@github-actions github-actions Bot changed the title [Triton/Gluon] # fp8 e4m3 RDNA3 dtype fix + tuned DSV4 A8W8 block-scale GEMM configs [Triton/Gluon] [JIT] # fp8 e4m3 RDNA3 dtype fix + tuned DSV4 A8W8 block-scale GEMM configs Sep 10, 2026
@github-actions github-actions Bot added the JIT label Sep 10, 2026
@amd-xavierwang amd-xavierwang changed the title [Triton/Gluon] [JIT] # fp8 e4m3 RDNA3 dtype fix + tuned DSV4 A8W8 block-scale GEMM configs [Triton/Gluon] [JIT] # RDNA# fp8 enablement + tuned DSV4 A8W8 block-scale GEMM configs Sep 10, 2026
@amd-xavierwang amd-xavierwang changed the title [Triton/Gluon] [JIT] # RDNA# fp8 enablement + tuned DSV4 A8W8 block-scale GEMM configs [Triton/Gluon] [JIT] # RDNA3 fp8 enablement + tuned DSV4 A8W8 block-scale GEMM configs Sep 10, 2026
Comment thread aiter/ops/triton/_triton_kernels/gemm/basic/gemm_a8w8_blockscale.py
Comment thread aiter/ops/triton/_triton_kernels/gemm/basic/gemm_a8w8_blockscale.py
Comment thread aiter/ops/triton/_triton_kernels/gemm/basic/gemm_a8w8_blockscale.py

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

@amd-xavierwang, can you please rebase on top of latest main branch? I think there are some CI failures in Aiter Test / Standard Tests suite that have already been fixed.

By this way, we can trigger CI again and then merge when everything is green.

@amd-xavierwang

Copy link
Copy Markdown
Contributor Author

CI looks good now!

@Boss2002n
Boss2002n merged commit a1d4231 into ROCm:main Sep 17, 2026
71 checks passed
@amd-xavierwang
amd-xavierwang deleted the gfx1100-blockscale-tuning-wip branch September 17, 2026 14:10
@brunomazzottiamd

Copy link
Copy Markdown
Contributor

CI looks good now!

I was trying to merge this in the last hour or so, but couldn't find an unresolved discussion... Now it's merged thanks to Satya!

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.

5 participants