Conversation
kashif
force-pushed
the
sync-fla-9d981ff
branch
2 times, most recently
from
September 7, 2026 13:42
cd74c4e to
92921d4
Compare
vasqu
reviewed
Sep 10, 2026
vasqu
left a comment
Collaborator
There was a problem hiding this comment.
Waitng on #1129 first so we can make sure we have tests for gdn and kda
I guess this sync would also be mainly for kda and gdn? I recall some support for context parallel, can we check in layers.py as well to sync the compatibility more concretely
kashif
force-pushed
the
sync-fla-9d981ff
branch
from
September 11, 2026 15:53
92921d4 to
4b11586
Compare
Update fla/torch-ext/fla/{ops,modules,utils} to fla-org/flash-linear-attention@516143e.
Adds new op families (atk, momentum_delta_rule, precond_gated_delta_rule, precond_kda,
and more) and picks up upstream fixes since the last sync.
- Rewrote new/changed absolute fla.* imports to relative, as before.
- Kept the minimal hand-maintained __init__.py/layers.py (no fla.layers/fla.models
vendoring); __init__.py now sets __version__, which utils/_decorators.py needs.
- Vendored pad_input/unpad_input into a new fla/torch-ext/fla/utils/padding.py so
fla.ops.deltaformer works without pulling in fla/layers.
- Bumped build.toml version 1 -> 2: chunk_gated_delta_rule/chunk_kda gained new
optional kwargs per check_public_api.py.
- Dropped triton_ascend backends everywhere (NPU-only, build.toml only declares
backends = ["cuda"]) and tilelang backends that were pure speed optimizations
(rwkv6, dplr, kda). Kept ops/common/backends/tilelang + its attn/common
registrations: it's the only fallback for a real Hopper correctness bug on
Triton 3.4.0-3.7.0 (see chunk_o.py's RuntimeError, upstream huggingface#640), not just perf.
- Also dropped the remaining scattered NPU-only code paths: ascend_compile_kwargs()
and the triton.language.extra.cann->ascend aliasing in utils/_compat.py, and the
Ascend grid-size workaround in modules/token_shift.py. Left the plain IS_NPU/IS_AMD/
IS_INTEL device-platform flags in utils/_device.py alone - they're a shared, cheap,
zero-dependency family, not NPU-specific code.
gdn2 gains dispatch() decorators on chunk_gdn2_bwd_wy_dqkg_fused/chunk_gdn2_fwd_intra
(upstream added a triton_ascend gdn2 backend to hang off them; not vendored here,
same as every other triton_ascend backend, and dispatch('gdn2') no-ops cleanly with
no registry present). solve_tril.py and utils/_device.py+__init__.py pick up AMD
gfx1250 TMA support (IS_AMD_TMA_ARCH, get_device_arch) - real, not NPU-specific.
Skipped ascend_ub_manager.py changes (file doesn't exist here) and the new
gdn2/backends/triton_ascend/ directory, consistent with prior NPU-code removal.
kashif
force-pushed
the
sync-fla-9d981ff
branch
from
September 17, 2026 07:57
4b11586 to
d2408df
Compare
check_kernel_test_versions.py expects tests to reference build.toml's declared version; bumped it to 2 there but missed the test file itself.
e52dbc0e..a7a32d8b: fixes dx/dw/db being written/counted for rows outside this program's tile when BT > BS (upstream #1249). Skipped the other two upstream commits in this window: one touches fla/layers/comba.py (not vendored), the other touches ops/kda/backends/triton_ascend (not vendored, same as always).
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.
Important: Read before submitting
Related issue
Closes #
What does this PR do?
Syncs the vendored
flakernel with the currentflash-linear-attentionmain (e52dbc0).Motivation
flawas a couple months stale. Routine refresh.Changes
fla/torch-ext/fla/{ops,modules,utils}from upstream, new absolute imports rewritten to relative as usual.deltaformerneedspad_input/unpad_inputfromfla/layers/utils.pyupstream (not vendored here) — originally pulled just those two functions intofla/torch-ext/fla/utils/padding.py; rebasing onto main picked up fla: fix relative import depths in the vendored source and add smoke tests #1129'sfla/torch-ext/fla/utils/layers.py(a full vendored copy, same idea, landed independently while this was open), so droppedpadding.pyand pointeddeltaformerat that instead.__init__.pynow sets__version__(upstream's_decorators.pyneeds it), kept everything else about it andlayers.pyas-is.build.tomlversion 1 → 2 (on top of main'sxpubackend addition),check_public_api.pyflagged new optional kwargs onchunk_gated_delta_rule/chunk_kda.triton_ascendbackends everywhere (NPU-only,build.tomldeclaresbackends = ["cuda", "xpu"], no NPU) andtilelangbackends that were pure speed optimizations (rwkv6, dplr, kda, and now gdn2 too — upstream added a triton_ascend gdn2 backend in the same window). Keptops/common/backends/tilelang(+ its attn/common registrations) since it's the only fallback for a real Hopper correctness bug on Triton 3.4.0-3.7.0, not just perf (see theRuntimeErrorinchunk_o.py, upstream metal-flash-sdpa: add or update card #640).ascend_compile_kwargs()and a tritonextra.cann→extra.ascendalias inutils/_compat.py, and an Ascend grid-size workaround inmodules/token_shift.py. Left the plainIS_NPU/IS_AMD/IS_INTELdevice-platform flags alone — shared, cheap, zero-dependency, not NPU-specific code.fla/tests/test_fla.py, both kept as-is) and fla: add XPU support and optimizations #1145 (XPU support — hand-merged itsIS_INTELadditions inchunk_o.py/chunk_delta_h.py/gated_delta_rule/{chunk_fwd,wy_fast}.pytogether with this branch's own upstream sync of those same files).flash-linear-attentionmovement since the last sync (516143e→e52dbc0):gdn2gained@dispatchdecorators (its new triton_ascend backend isn't vendored, same as everywhere else —dispatch()no-ops cleanly with no registry present), andsolve_tril.py/utils/_device.py/utils/__init__.pypicked up real AMD gfx1250 TMA support.Testing
flash-linear-attentionitself): a missing/mismatchedtilelanginstall, and a real upstream bug inops/delta_rule/parallel.py(imports a namewy_fast.pydoesn't define).check_public_api.pyafter every change — same two additive kwarg changes throughout, nothing new.Checklist
🤖 Generated with Claude Code