Skip to content

[Perf] TileLang backends for linear-attention chunk backward (simple_gla / gla / delta_rule / gdn) - #1046

Open
markovchain-builder wants to merge 8 commits into
fla-org:mainfrom
markovchain-builder:argus/tilelang-fla
Open

[Perf] TileLang backends for linear-attention chunk backward (simple_gla / gla / delta_rule / gdn)#1046
markovchain-builder wants to merge 8 commits into
fla-org:mainfrom
markovchain-builder:argus/tilelang-fla

Conversation

@markovchain-builder

Copy link
Copy Markdown
Contributor

Summary

Hand-written TileLang backends for the backward path of four core linear-attention chunk kernels, gated by FLA_TILELANG with automatic Triton fallback. All four pass the frozen fwd+bwd gate (NaN-poison) and show measured speedups over the existing Triton kernels on H100 NVL.

Results (N=3 geomean, fwd+bwd)

Op fwdbwd geomean all-row peak shape baseline
chunk_gla +20.0% +9.2% 1.32× pure Triton (FLA_TILELANG=0 vs =1)
chunk_simple_gla +16.1% +8.0% 1.22× pure Triton
chunk_delta_rule +5.6% +4.4% 1.155× (D256) verify --base main
chunk_gdn +3.7% +2.1% verify --base main

How

All speedups come from the backward path — fusing the separate Triton backward kernels (dq/dk/dw/dg + WY-representation backward) into single producer–consumer TileLang pipelines that keep intermediate state on-chip and cut HBM round trips:

  • common: shared TileLang chunk-backward backend (dqkwg staging, chunk_h) + a CUDA 12.6/12.7 __nv_fp8_e8m0 stub so TileLang's FP8 header compiles on this toolchain.
  • simple_gla: V-first direct-state dqkwg, K-inner staging, shadow-state reuse, fused hdh epilogue.
  • gla: op-local fused K-tile backward consumer (dense D128/D256).
  • delta_rule: D256-focused fused producer–consumer WY-backward route.
  • gdn: split rewrite of the prepare_wy backward (compiled via CUDA 12.8 for Hopper FP8 e8m0 + TMA).

Methodology

  • Correctness: the repo's frozen tests/ops/test_*.py (fwd+bwd, NaN memory poisoning), unchanged; naive references unchanged.
  • Timing: benchmarks/ops/run.py / verify.py, N=3, geomean, identical shape/precision/numeric flags on both sides.
  • Hardware: NVIDIA H100 NVL · PyTorch 2.7.1+cu126 · Triton 3.3.1 · TileLang 0.1.9.
  • Backend is opt-in via FLA_TILELANG=1 and falls back to Triton when no usable nvcc is present.

argus added 5 commits July 19, 2026 13:36
Shared hand-written TileLang backend for FLA linear-attention chunk backward
(dqkwg staging + chunk_h), wired into the backend dispatch. Includes a
CUDA 12.6/12.7 __nv_fp8_e8m0 stub so TileLang's FP8 header compiles on this
toolchain. Backend is gated by FLA_TILELANG and falls back to Triton when
TileLang has no usable nvcc.
V-first direct-state dqkwg with K-inner staging and on-chip shadow-state reuse,
fusing the hdh epilogue to cut HBM round trips. Frozen N=3 scorer, pure-Triton
baseline (FLA_TILELANG=0 vs 1) on H100 NVL: fwdbwd geomean +16.1%, all-row
+8.0%, peak 1.22x; fwd+bwd NaN-poison gate green.
Op-local fused K-tile backward consumer for dense D128/D256, keeping the
existing dh/dv/dA producers. Frozen N=3 scorer, pure-Triton baseline on
H100 NVL: fwdbwd geomean +20.0%, all-row +9.2%, peak 1.32x; gate green.
D256-focused fused producer-consumer WY-backward route, collapsing the WY
representation backward into one on-chip pipeline to remove dhu / WY-backward
HBM round trips. Frozen N=3 scorer (verify --base main): fwdbwd geomean +5.6%,
all-row +4.4%, D256 peak 1.155x; gate green.
Split rewrite of the gated delta net (GDN) prepare_wy backward, compiled via
the CUDA 12.8 toolkit (Hopper FP8 e8m0 + TMA). Frozen N=3 scorer: fwdbwd
geomean +3.7%, all-row +2.1%, three repeats all GATE PASSED; correctness green.
@zhiyuan1i

Copy link
Copy Markdown
Collaborator

Thanks to the all-powerful agents. This MR is simply enormous…

@zhiyuan1i

Copy link
Copy Markdown
Collaborator

Please restore the tl.debug_barrier() in prepare_wy_repr_bwd_kernel (fla/ops/gated_delta_rule/wy_fast.py) — this PR removes it, but the barrier is load-bearing.

It was added deliberately in #944: without it, this kernel misbehaves on the triton-ascend backend (a timing/synchronization issue in that Triton variant — the barrier orders the preceding b_dA dot/where sequence before the K-loop that recomputes b_A and performs the p_dk load-modify-store). It may look redundant on the NVIDIA path, but dropping it unconditionally re-exposes the race on the affected backend and fails there silently.

If you have evidence the underlying backend issue is fixed (a specific Triton/backend changelog), please gate the removal on a minimum-version check instead of deleting the barrier outright.

The barrier was added deliberately in fla-org#944: without it, triton-ascend can
misorder the b_dA update before the dk load-modify-store in the K-loop,
failing silently on that backend. Restore it with a comment so it is not
removed again as an NVIDIA-only-looking no-op.

Patch from @lbx154, applied via maintainer edit.
@chen-xiao-yu-1

Copy link
Copy Markdown

Follow-up to #1046: restored tl.debug_barrier() in prepare_wy_repr_bwd_kernel, preserving the synchronization barrier introduced in #944.

Fix: chen-xiao-yu-1/flash-linear-attention@2a981b71

@chen-xiao-yu-1

Copy link
Copy Markdown

Follow-up on the merge conflict in #1046:

I resolved the current conflict locally and pushed the prepared branch here:

chen-xiao-yu-1:pr-1046-conflict-resolved
The branch contains:

4906a2e — merge latest origin/main into the PR branch and resolve the conflict
dfc334d — fix optional TileLang import for delta rule
The conflict was in fla/ops/common/chunk_h.py.
The resolution keeps the explicit pointer+mask boundary handling from main while preserving this PR's STORE_MMA_STATE / FUSE_HDH_LAST logic.
Additional issue found while testing:
In environments without tilelang, importing fla.ops can fail even with FLA_TILELANG=0, because fla/ops/delta_rule/chunk.py imported the TileLang backend at module import time. I changed that to a lazy import inside the _can_use_tilelang_wy_dqkw_fused(...) guarded path.

@chen-xiao-yu-1

Copy link
Copy Markdown

The current merge conflict is resolved in markovchain-builder#1.

Resolve chunk_h.py: main migrated off tl.make_block_ptr to raw pointer
arithmetic; keep the STORE_MMA_STATE dual-write and FUSE_HDH_LAST paths,
rewritten with the same raw-pointer + mask style.

@zhiyuan1i zhiyuan1i left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

Thanks for your patience on this one. I've gone ahead and pushed two things to the branch to help move it along:

  • Resolved the textual merge conflict in fla/ops/common/chunk_h.py (main's raw-pointer migration × your STORE_MMA_STATE dual-write and FUSE_HDH_LAST paths, rewritten in the raw-pointer + mask style; cross-checked against @chen-xiao-yu-1's resolution — semantically equivalent).
  • Restored the tl.debug_barrier() in prepare_wy_repr_bwd_kernel with @lbx154's patch from the comments (load-bearing on triton-ascend, now with a why-comment).

However, CI on the merged head shows this PR now needs a rebase-level adaptation to main, which is beyond a conflict fix and really needs the author:

  1. test-ops fails: the NPU call sequence in fla/ops/gated_delta_rule/wy_fast.py (the per-kernel prepare_wy_repr_bwd_k_npu / _v_npu / _da_mask_npu / ... launches) no longer matches main's triton_ascend backend, which was restructured into a fused prepare_wy_repr_bwd_kv_npu with different parameters and grid semantics (task_num/num_core persistence, G_EXP_PRECOMP, etc.). The test failure is AttributeError: ... no attribute 'prepare_wy_repr_bwd_k_npu'. The dispatch sequence needs to be rewritten against the new fused backend (or dropped in favor of the existing backend dispatch).

  2. lint fails: main added a no-make-block-ptr pre-commit hook (kernels migrated off tl.make_block_ptr/tl.advance for triton-main compatibility). fla/ops/delta_rule/backends/triton/chunk_bwd.py on this branch still has 19 block_ptr uses that need migrating to the raw-pointer + mask style.

  3. check-pr-title fails: main now requires the PR body checklist to be fully ticked. Please refresh the body against the current template.

Could you rebase the branch onto current main and address these three? Happy to re-review quickly once CI is green.

@chen-xiao-yu-1

Copy link
Copy Markdown

@zhiyuan1i I prepared the requested fixes on chen-xiao-yu-1:pr-1046-conflict-resolved.

Could you merge current main into the PR branch and cherry-pick these two commits?

  • bf8aa1c3 — retain the lazy TileLang import fix
  • 7d6304c7 — use main's existing fused NPU backend dispatch and migrate the 19 block-pointer accesses to raw pointers and masks

The prepared branch preserves the reviewed chunk_h.py resolution and the tl.debug_barrier() why-comment.

The PR body checklist also needs to be refreshed against the current template.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

waiting-author Reviewer acted; ball is in the author's court

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants