Fuse block-banking cat into attn_res_gate via close_block/WRITE_BLOCK… - #4726
Open
yanxuer-999 wants to merge 2 commits into
Open
Fuse block-banking cat into attn_res_gate via close_block/WRITE_BLOCK…#4726yanxuer-999 wants to merge 2 commits into
yanxuer-999 wants to merge 2 commits into
Conversation
…_CAT Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
🏷️ CI GuideRuns automatically on every PR:
Extended tests (opt-in via labels):
|
- Skip redundant reshape()/contiguous() dispatcher calls in attn_res_gate when q/prefix/block_residual/add_hidden* are already the exact shape/layout the kernel wants (the common case for ATOM's AttnRes, whose inputs are already contiguous from the previous layer). - Guard the per-call ATTN_RES_GATE debug log behind isEnabledFor(INFO) so f-string formatting isn't paid on every call when info logging is disabled. - Add ATTN_RES_TRITON_AUTOTUNE (default off), an fla-style live-search escape hatch that wraps attnres_fwd_kernel with @triton.autotune over (BL, num_warps, num_stages). Production still uses the static per-token-count table in the wrapper (bounded compile cost, CUDAGraph-safe); this flag exists to regenerate/verify that table offline. Confirmed the table is already at/near the tiling optimum for Kimi-K3's real (T, B) shapes. Co-authored-by: Cursor <cursoragent@cursor.com>
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
1 participant
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.
…_CAT
Motivation
Technical Details
Test Plan
Test Result
Submission Checklist