[CK_TILE] FMHA BWD Use Persistent Kernels in Deterministic Mode - #2216
Merged
Conversation
Contributor
🏷️ CI GuideRuns automatically on every PR:
Extended tests (opt-in via labels):
|
Contributor
There was a problem hiding this comment.
Pull request overview
Enable/use CK Tile FMHA backward persistent-kernel path in deterministic mode by deriving dq_accum split count (nsplits) from the launcher traits and simplifying accumulator allocation.
Changes:
- Compute FMHA BWD traits earlier and use launcher-provided
dq_acc_splitsto sizedq_accum. - Unify
dq_accumallocation to always usensplits(instead of separate deterministic/non-deterministic branches). - Add
-fbracket-depth=1024to HIP compile flags; bumpcomposable_kernelsubmodule.
Reviewed changes
Copilot reviewed 4 out of 4 changed files in this pull request and generated 5 comments.
| File | Description |
|---|---|
| csrc/py_itfs_ck/mha_varlen_bwd_kernels.cu | Derive nsplits from FMHA BWD launcher traits and allocate dq_accum accordingly for varlen BWD. |
| csrc/py_itfs_ck/mha_bwd_kernels.cu | Same nsplits/dq_accum refactor for non-varlen BWD; moves bias/traits setup earlier. |
| aiter/jit/optCompilerConfig.json | Adds extra HIP compiler flag (-fbracket-depth=1024) in multiple configs. |
| 3rdparty/composable_kernel | Updates the CK submodule revision. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
DDEle
force-pushed
the
ck-fmha-bwd-persistent-det
branch
2 times, most recently
from
March 12, 2026 09:53
7404493 to
f330c92
Compare
DDEle
marked this pull request as ready for review
March 13, 2026 01:26
DDEle
marked this pull request as draft
March 13, 2026 01:27
DDEle
force-pushed
the
ck-fmha-bwd-persistent-det
branch
from
March 13, 2026 03:40
f330c92 to
06d29f3
Compare
DDEle
force-pushed
the
ck-fmha-bwd-persistent-det
branch
from
March 13, 2026 06:33
06d29f3 to
7c4242f
Compare
DDEle
marked this pull request as ready for review
March 13, 2026 07:19
AMD-yanfeiwang
pushed a commit
to AMD-yanfeiwang/aiter
that referenced
this pull request
Mar 18, 2026
DDEle
added a commit
that referenced
this pull request
Apr 28, 2026
PR #2321 inlined launcher.dq_acc_splits / needs_zero_dq_acc as a hardcoded `nsplits = ceil(seqlen_k/16)` and unconditional zero, because the CK branch it bumped to had temporarily removed fmha_bwd_launcher. The pinned CK now has the launcher back, so restore the #2216 pattern: construct fmha_bwd_traits + fmha_bwd_launcher and read nsplits and needs_zero_dq_acc from it. Functional behavior is unchanged with the current pinned CK; this is a prep commit to keep the diff for the upcoming #6152 (unified workspace) adaptation small. Touched files: - csrc/py_itfs_ck/mha_bwd_kernels.cu (batch CK entry) - csrc/py_itfs_ck/mha_varlen_bwd_kernels.cu (group CK entry) - op_tests/cpp/mha/benchmark_mha_bwd.cpp (benchmark host)
valarLip
pushed a commit
that referenced
this pull request
May 14, 2026
* [CK_TILE] mha bwd: switch to fmha_bwd_launcher usage PR #2321 inlined launcher.dq_acc_splits / needs_zero_dq_acc as a hardcoded `nsplits = ceil(seqlen_k/16)` and unconditional zero, because the CK branch it bumped to had temporarily removed fmha_bwd_launcher. The pinned CK now has the launcher back, so restore the #2216 pattern: construct fmha_bwd_traits + fmha_bwd_launcher and read nsplits and needs_zero_dq_acc from it. Functional behavior is unchanged with the current pinned CK; this is a prep commit to keep the diff for the upcoming #6152 (unified workspace) adaptation small. Touched files: - csrc/py_itfs_ck/mha_bwd_kernels.cu (batch CK entry) - csrc/py_itfs_ck/mha_varlen_bwd_kernels.cu (group CK entry) - op_tests/cpp/mha/benchmark_mha_bwd.cpp (benchmark host) * [CK_TILE] mha bwd: adapt to CK #6152 unified workspace API CK PR #6152 replaces fmha_bwd_args.dq_acc_ptr + the four dq_acc_* strides with a single opaque workspace_ptr, and exposes fmha_bwd_launcher.workspace_size + prepare_workspace(void*) instead of dq_acc_splits + needs_zero_dq_acc. aiter::mha_bwd_args is unchanged (no new fields). The existing dq_acc_ptr field doubles as the CK workspace pointer for the CK path; the four dq_acc_* stride fields are kept for the ASM v3 path that still consumes them. The torch entries (mha_bwd_kernels.cu / mha_varlen_bwd_kernels.cu / benchmark_mha_bwd.cpp) construct the launcher, allocate an at::Tensor workspace of launcher.workspace_size bytes, call launcher.prepare_workspace, and pass workspace.data_ptr() through dq_acc_ptr. For group mode, the launcher requires host-side seqstart arrays. The torch varlen entry copies cu_seqlens_q_padded (or cu_seqlens_q if no padded variant was provided) via .to(at::kCPU) — using the same "physical seqstart" convention that the kernel itself indexes dq_acc with. The aiter::mha_bwd dispatcher does its own hipMemcpy from a.seqstart_q_ptr / a.seqstart_k_ptr for the CK fallback path. Neither host buffer crosses the call boundary. Submodule bumped to 8a59f8afa58 (subtree-split of monorepo users/yiding12/fmha-bwd-workspace tip, includes the per-nhead dq_acc stride fix for group mode). * [CK_TILE] mha bwd: single workspace_alloc callback for both paths Removes the dual-purpose dq_acc_ptr field and the four dq_acc_* stride fields from aiter::mha_bwd_args, replacing them with one callback that serves both dispatch paths: std::function<void*(size_t bytes, bool zero_init)> workspace_alloc; - CK fallback (in aiter::mha_bwd): Constructs fmha_bwd_launcher, queries workspace_size, calls workspace_alloc(size, zero_init=false), then forwards the pointer to launcher.prepare_workspace + launcher.run. The torch entries no longer construct fmha_bwd_launcher themselves; py_itfs_ck/* shrinks considerably (traits + launcher + workspace blocks deleted, varlen also drops its private cu_seqlens D2H copy). - ASM v3 (in fmha_v3_bwd): Determines dq_accum shape and byte count internally (based on is_group_mode, v3_atomic_fp32, hdim, batch, nhead, seqlen) and calls workspace_alloc(bytes, zero_init=true). The torch entries no longer allocate or zero dq_accum themselves; py_itfs_cu/asm_mha_*.cu shrinks to a small lambda. The zero_init flag lets each backing storage pick its efficient zero path (torch::zeros, DeviceMem::SetZero, ...) instead of forcing the dispatch path to know HIP memset semantics. ASM kernels need it because they atomically accumulate into dq_accum; the CK launcher fills its own workspace and does not need a pre-zero. * [CK_TILE] mha bwd: address review comments on workspace_alloc - mha_bwd CK fallback: explicitly reject group mode with missing seqstart pointers (LOG_ERROR + return -1) instead of silently passing nullptr to the launcher. Matches the existing AITER_LOG_WARNING pattern in fmha_v3_bwd for unsupported configurations. - benchmark workspace_alloc: replace silent return-nullptr-on-oversize with AITER_CHECK so the failure is loud at the actual call site rather than as an opaque kernel crash later. Also tighten the zero-init path to hipMemset only `bytes` rather than the full pre-allocated buffer. * Update CK pin as ROCm/rocm-libraries#6152 merged
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
This PR, together with ROCm/rocm-libraries#5174, enables a persistent-kernel execution path for FMHA backward (dQ/dK/dV) in deterministic mode, adjusting how dQ accumulation is split, stored, and converted back to final gradients.
Technical Details
Test Plan
CK hosted AITER tests Before ROCm/rocm-libraries#5174 merged: http://micimaster.amd.com/blue/organizations/jenkins/rocm-libraries-folder%2FComposable%20Kernel/detail/PR-5174/7/pipeline/207
Test Result
Submission Checklist