Skip to content

[CI][CK_TILE] Update CK and fix fmha_fwd arg init#2334

Open
hyoon1 wants to merge 1 commit intomainfrom
fix/nightly-ci-ck-fmha-fwd
Open

[CI][CK_TILE] Update CK and fix fmha_fwd arg init#2334
hyoon1 wants to merge 1 commit intomainfrom
fix/nightly-ci-ck-fmha-fwd

Conversation

@hyoon1
Copy link

@hyoon1 hyoon1 commented Mar 18, 2026

Motivation

Nightly CI failed at compile-time in mha_fwd.cu with the latest CK.

Technical Details

CK_TILE fmha_fwd_args added new fields (num_head_q_total, head_start).
Insert two arguments in order to match updated struct layout.

  • Update 3rdparty/composable_kernel submodule to e5683e22902109f7652d2a1c3d39ed4c525f90ef
  • Fix FMHA forward arg initialization in csrc/cpp_itfs/mha_fwd.cu

Test Plan

Test Result

Submission Checklist

@hyoon1 hyoon1 requested review from a team, Copilot, illsilin and rocking5566 March 18, 2026 17:26
@github-actions
Copy link
Contributor

🏷️ CI Guide

Runs automatically on every PR:

  • ✅ Pre-checks (submodule verification, code formatting)
  • ✅ Aiter op tests (gfx942 + gfx950)
  • ✅ Triton tests (only when aiter/ops/triton/** or related paths are changed)

Extended tests (opt-in via labels):

Label Tests
ci:sglang SGLang integration tests
ci:atom ATOM benchmark (DeepSeek-R1 + GPT-OSS)
ci:vllm vLLM benchmark
ci:all All of the above

Add labels via the sidebar or gh pr edit 2334 --add-label <label>

Copy link
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

Updates the Composable Kernel (CK) submodule and adjusts FMHA forward argument initialization to match CK’s updated fmha_fwd_args layout so Nightly CI compiles successfully.

Changes:

  • Bump 3rdparty/composable_kernel submodule to commit e5683e2290....
  • Add the two newly required fmha_fwd_args fields (num_head_q_total, head_start) to the FMHA forward argument list.

Reviewed changes

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

File Description
csrc/cpp_itfs/mha_fwd.cu Adds two arguments to align with CK’s updated fmha_fwd_args struct layout.
3rdparty/composable_kernel Updates CK submodule revision to the version that introduced the struct layout change.

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

You can also share your feedback on Copilot code review. Take the survey.

Comment on lines 315 to 322
a.hdim_v,
a.nhead_q,
a.nhead_k,
0, // num_head_q_total
0, // head_start
a.scale_s,
a.logits_soft_cap,
a.stride_q,
Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

For non-head-sliced paths, keeping these fields as 0 is fine. CK treats them as optional and uses internal fallback behavior for the default (non-sliced) case.

@hyoon1 hyoon1 force-pushed the fix/nightly-ci-ck-fmha-fwd branch from af75cb8 to ba1680c Compare March 18, 2026 17:30
@hyoon1 hyoon1 changed the title [CI][Nightly] Update CK and fix fmha_fwd arg init [CI][CK_TILE] Update CK and fix fmha_fwd arg init Mar 18, 2026
Copy link
Contributor

@rocking5566 rocking5566 left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM

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.

3 participants