Skip to content

Jim/dev/mi400 mha bwd#3281

Open
slippedJim wants to merge 4 commits into
mainfrom
jim/dev/mi400_mha_bwd
Open

Jim/dev/mi400 mha bwd#3281
slippedJim wants to merge 4 commits into
mainfrom
jim/dev/mi400_mha_bwd

Conversation

@slippedJim
Copy link
Copy Markdown
Contributor

Motivation

Technical Details

Test Plan

Test Result

Submission Checklist

@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: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 3281 --add-label <label>

Copy link
Copy Markdown
Contributor

Copilot AI left a comment

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 adds FMHA v3 backward (MHA bwd) ASM-kernel support for the gfx1250 GPU architecture by introducing new gfx1250 kernel/config entries and updating the runtime dispatch logic to recognize and launch them.

Changes:

  • Add gfx1250 FMHA v3 bwd config CSVs for ODO (pre), dQ/dK/dV, and dQ convert (post) kernels.
  • Extend fmha_v3_bwd dispatch to allow gfx1250, select the appropriate post-processing config, and always enable post-processing on gfx1250.
  • Adjust workgroup size (bdx) to 128 on gfx1250 for pre/dqdkdv/post kernel launches.

Reviewed changes

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

File Description
hsa/gfx1250/fmha_v3_bwd/fmha_bwd_odo.csv Adds gfx1250 ODO (pre-processing) kernel configuration entry.
hsa/gfx1250/fmha_v3_bwd/fmha_bwd_dqdkdv.csv Adds gfx1250 dQ/dK/dV kernel configuration entries (including causal BR).
hsa/gfx1250/fmha_v3_bwd/fmha_bwd_dq_convert.csv Adds gfx1250 dQ convert (post-processing) kernel configuration entry.
csrc/cpp_itfs/mha_bwd.cu Enables gfx1250 in asm v3 bwd dispatch, sets post-processing behavior, and tweaks launch block size.

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

Comment thread csrc/cpp_itfs/mha_bwd.cu
Comment on lines 314 to 326
@@ -321,11 +321,13 @@ float fmha_v3_bwd(mha_bwd_args a, const ck_tile::stream_config& s)
{
return static_cast<CFG*>(nullptr);
}
} else {
return &cfg_fmha_bwd_dq_convert; // gfx1250 only support atomic32=1
}
Comment thread csrc/cpp_itfs/mha_bwd.cu

if((mt == 1) || (mt == 2))
{ // causal
{ // mask kb
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants