Skip to content

fla: sync with flash-linear-attention main - #1141

Open
kashif wants to merge 5 commits into
huggingface:mainfrom
kashif:sync-fla-9d981ff
Open

kashif wants to merge 5 commits into
huggingface:mainfrom
kashif:sync-fla-9d981ff

Conversation

@kashif

@kashif kashif commented Sep 7, 2026

Copy link
Copy Markdown
Contributor

Important: Read before submitting

New contributors: Please open or comment on an issue before submitting a
PR to discuss the change you'd like to make. This helps us align on approach and
avoids wasted effort on changes we may not be able to merge. Please only create
a PR once one of the project maintainers agrees on your outlined approach.

PRs of contributors who are not vouched for are automatically closed. Regular
contributors are added to the vouch list.

The kernels-community repository is for:

  1. Kernels developed by Hugging Face and partners.
  2. Kernels developed by third parties, that have not been 'kernelized' yet, but
    are used by Hugging Face projects such as diffusers and transformers.

We cannot accept PRs for the following:

  • New kernels that are not in categories (1) and (2).
  • Changes to kernels in category (2), these should be submitted upstream.

For, LLM-generated changes, we prefer that you write your prompt in an issue
over a PR with LLM-generated changes.

Related issue

Closes #

What does this PR do?

Syncs the vendored fla kernel with the current flash-linear-attention main (e52dbc0).

Motivation

fla was a couple months stale. Routine refresh.

Changes

  • Updated fla/torch-ext/fla/{ops,modules,utils} from upstream, new absolute imports rewritten to relative as usual.
  • New op families came along: atk, momentum_delta_rule, precond_gated_delta_rule, precond_kda, and a few more.
  • deltaformer needs pad_input/unpad_input from fla/layers/utils.py upstream (not vendored here) — originally pulled just those two functions into fla/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's fla/torch-ext/fla/utils/layers.py (a full vendored copy, same idea, landed independently while this was open), so dropped padding.py and pointed deltaformer at that instead.
  • __init__.py now sets __version__ (upstream's _decorators.py needs it), kept everything else about it and layers.py as-is.
  • Bumped build.toml version 1 → 2 (on top of main's xpu backend addition), check_public_api.py flagged new optional kwargs on chunk_gated_delta_rule/chunk_kda.
  • Dropped triton_ascend backends everywhere (NPU-only, build.toml declares backends = ["cuda", "xpu"], no NPU) and tilelang backends that were pure speed optimizations (rwkv6, dplr, kda, and now gdn2 too — upstream added a triton_ascend gdn2 backend in the same window). Kept ops/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 the RuntimeError in chunk_o.py, upstream metal-flash-sdpa: add or update card #640).
  • Also swept the remaining scattered NPU-only code: ascend_compile_kwargs() and a triton extra.cannextra.ascend alias in utils/_compat.py, and an Ascend grid-size workaround in modules/token_shift.py. Left the plain IS_NPU/IS_AMD/IS_INTEL device-platform flags alone — shared, cheap, zero-dependency, not NPU-specific code.
  • Rebased onto main to pick up fla: fix relative import depths in the vendored source and add smoke tests #1129 (relative-import fixes + fla/tests/test_fla.py, both kept as-is) and fla: add XPU support and optimizations #1145 (XPU support — hand-merged its IS_INTEL additions in chunk_o.py/chunk_delta_h.py/gated_delta_rule/{chunk_fwd,wy_fast}.py together with this branch's own upstream sync of those same files).
  • Picked up the small amount of upstream flash-linear-attention movement since the last sync (516143ee52dbc0): gdn2 gained @dispatch decorators (its new triton_ascend backend isn't vendored, same as everywhere else — dispatch() no-ops cleanly with no registry present), and solve_tril.py/utils/_device.py/utils/__init__.py picked up real AMD gfx1250 TMA support.

Testing

  • Imported every submodule of the built package. Everything imports cleanly except two pre-existing issues, neither introduced here (reproduced identically against flash-linear-attention itself): a missing/mismatched tilelang install, and a real upstream bug in ops/delta_rule/parallel.py (imports a name wy_fast.py doesn't define).
  • Re-ran check_public_api.py after every change — same two additive kwarg changes throughout, nothing new.

Checklist

  • This PR is linked to an issue that was discussed and approved
  • I have tested these changes locally
  • New/changed functionality has test coverage
  • LLM disclosure:
    • I did not use an LLM to create this PR.
    • I used and LLM for assistance while creating this PR.
    • This PR was mostly or completely generated by an LLM.

🤖 Generated with Claude Code

@github-actions github-actions Bot added the chore Version bumps, releases, misc maintenance label Sep 7, 2026
@kashif
kashif force-pushed the sync-fla-9d981ff branch 2 times, most recently from cd74c4e to 92921d4 Compare September 7, 2026 13:42

@vasqu vasqu 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.

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

@github-actions github-actions Bot added the needs-rebase Has merge conflicts with the base branch label Sep 11, 2026
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.
@github-actions github-actions Bot removed the needs-rebase Has merge conflicts with the base branch label Sep 17, 2026
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).
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

chore Version bumps, releases, misc maintenance

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants