name = flash-linear-attention (fla-core)
version = 0.5.1
tag = v0.5.1
upstream = https://github.com/fla-org/flash-linear-attention
pypi = https://pypi.org/project/fla-core/0.5.1/
license = MIT
license_file = LICENSE
copyright = 2023-2026 Songlin Yang, Yu Zhang, Zhiyuan Li

description =
    Minimal, pinned, pruned snapshot of fla-core 0.5.1. Only the transitive file
    closure needed by the gated-delta-rule fast path is vendored, so Unsloth can
    enable the Triton kernels for the Qwen3.5 / Qwen3.6 / Qwen3-Next
    gated-deltanet models, the GLM-5.3-Flash / Kimi-Linear Kimi delta
    attention models and the Ling 2.5 / 2.6 / 3.0 remote-code models without
    requiring `pip install flash-linear-attention`.

exported_api =
    fla.modules.FusedRMSNormGated
    fla.modules.ShortConvolution
    fla.ops.gated_delta_rule.chunk_gated_delta_rule
    fla.ops.gated_delta_rule.fused_recurrent_gated_delta_rule
    fla.ops.kda.chunk_kda
    fla.ops.kda.fused_recurrent_kda
    fla.ops.simple_gla.chunk.chunk_simple_gla
    fla.ops.simple_gla.fused_recurrent.fused_recurrent_simple_gla

modifications =
    - Narrowed package __init__ exports for fla, fla.ops, fla.modules and
      fla.ops.gated_delta_rule so importing the package does not eagerly pull
      fla.layers / fla.models / "import every op" (see the "Modified by Unsloth"
      note atop each narrowed __init__.py). All other files are verbatim upstream
      with their original MIT headers preserved, except the backported fixes below.
    - Backported three post-v0.5.1 upstream correctness fixes (each marked with an
      inline "Unsloth: backported from fla PR #..." comment):
        * PR #953 (issue #945) in ops/common/chunk_delta_h.py: pin the Blackwell
          fwd-h kernel to num_warps=2 (the 4-warp config hits a Triton tl.dot
          recurrence race that silently corrupts h / v_new on B200).
        * PR #1000 (issue #999) in ops/gated_delta_rule/wy_fast.py: restrict the
          Blackwell prepare_wy_repr_bwd_kernel autotune to the B200-validated
          config (unstable configs can hang / misaligned-address the bwd).
          Narrowed by Unsloth: PR #1000 pins num_warps=2 AND num_stages=4, but only
          num_warps was implicated by #999 and the report was on triton 3.3.1, while
          num_stages=4 is the expensive half. Measured on one B200 (torch 2.12.1+cu130,
          triton 3.7.1), chunk_gated_delta_rule fwd+bwd B=2 T=8192 H=48 D=128 bf16:
          the full pin costs 0.733 ms of 5.309 ms and num_stages alone accounts for
          0.624 ms of that. So the warp pin is kept unconditionally and num_stages is
          allowed to autotune from triton 3.6 onward; below 3.6 the exact upstream pin
          is kept, since that is the era #999 was reported against.
        * PR #983 (issue #640) in ops/common/chunk_o.py (+ TRITON_ABOVE_3_7_1 in
          utils/_compat.py and utils/__init__.py): narrow the Hopper gated
          chunk_bwd_dqkwg guard to Triton [3.4.0, 3.7.1); 3.7.1 fixes the bug.
          Unsloth goes one step further than upstream here and does not merely
          narrow the guard, it removes the need for it. Issue #640's own root-cause
          bisection pins the miscompile to BK == 64 on Hopper (BK 32 and BK 128 both
          measured clean; the autotune config space was explicitly ruled out), and
          Hopper's CONST_TILING of 128 only reaches BK 64 for head dims 33..64. So
          chunk_bwd_dqkwg steps those down to BK 32 and keeps the Triton fast path,
          and the RuntimeError is retained only as an unreachable fence. Its
          remediation text was also rewritten: upstream points at
          `pip install tilelang`, which cannot work here because the TileLang
          kernels are pruned (below) and _inject_vendored_fla force-sets
          FLA_TILELANG=0.
          The Hopper test is also widened from upstream's ``IS_NVIDIA_HOPPER``
          alone, which fla/utils/_device.py freezes at import from device 0, to
          that OR a probe of ``k.device`` (``_device_is_nvidia_hopper``): the
          launcher already picks CONST_TILING from ``k.device.index``, so on a
          mixed host a Hopper card at a nonzero index would otherwise keep the
          miscompiled tile. The probe uses compute capability 9, not
          ``check_shared_mem('hopper')``, which is a >=232448-byte shared-memory
          tier test that Blackwell B200 also passes.
    - Added the fla.ops.simple_gla closure for remote-code lightning-attention
      models that import it directly (inclusionAI Ling 2.5 / 2.6,
      BailingMoeV2_5: `from fla.ops.simple_gla.chunk import chunk_simple_gla`,
      `from fla.ops.simple_gla.fused_recurrent import fused_recurrent_simple_gla`):
      ops/simple_gla/{chunk,fused_recurrent}.py and ops/common/{chunk_h,
      fused_recurrent}.py verbatim from fla-core 0.5.1. ops/simple_gla/__init__.py
      is narrowed to those two entry points (upstream also imports fused_chunk and
      parallel, whose closures are not vendored). The backward reuses the vendored
      ops/common/chunk_o.py, so it keeps the Hopper chunk_bwd_dqkwg fix above.
    - Added the fla.ops.kda closure for the Kimi delta attention models
      (transformers glm5_next and kimi_linear resolve `chunk_kda` /
      `fused_recurrent_kda` from fla through their kernel-hub wrappers):
      ops/kda/{__init__,chunk,chunk_bwd,chunk_fwd,chunk_intra,
      chunk_intra_token_parallel,fused_recurrent,gate,wy_fast}.py,
      ops/kda/backends/flashkda.py and ops/gla/{__init__,chunk,fused_chunk,
      fused_recurrent}.py (chunk_kda reuses chunk_gla_fwd_o_gk) verbatim from
      fla-core 0.5.1. ops/kda/backends/__init__.py is narrowed to register only
      the FlashKDA backend (optional `flash_kda` package, inference only); the
      TileLang KDA backend and its kernel backends/tilelang/chunk_bwd_dqkg.py
      are not vendored. ops/kda/naive.py and ops/gla/naive.py are dropped
      (reference implementations; the first needs einops). The KDA files share
      the vendored ops/common/chunk_delta_h.py, ops/cp and ops/utils, so they
      run on the #953 fix above.
    - Added the fla.modules ShortConvolution closure for Ling 3.0
      (BailingMoeV3) remote code, which imports it beside the KDA ops
      (`from fla.modules import FusedRMSNormGated, ShortConvolution`):
      modules/conv/{short_conv,causal_conv1d}.py, modules/conv/triton/{__init__,
      kernels,ops}.py, modules/conv/cp/{__init__,ops}.py and
      modules/conv/cuda/{__init__,ops}.py verbatim from fla-core 0.5.1, plus
      two upstream fixes in short_conv.py: fla PR #1257 (FLA_CONV_BACKEND=cuda
      without causal-conv1d installed now falls back to Triton instead of
      calling the missing function) and the short_conv hunk of fla PR #1203
      (a packed batch takes the single-token decode shortcut only when every
      sequence has exactly one token).
      modules/conv/__init__.py is narrowed to ShortConvolution and
      causal_conv1d (long_conv.py is not vendored), and modules/backends/
      __init__.py registers no backend, since the Triton-Ascend NPU kernels
      under modules/backends/triton_ascend are not vendored and the tree is
      only injected on CUDA. fla.modules resolves ShortConvolution lazily: its
      files import einops, which the gated-deltanet path does not need. Later
      upstream conv fixes are not backported, as none reaches Ling 3.0's
      training or decode path: #1097 (more than 65535 sequences per decode
      launch) and #1139 / #1141 (gradients into a conv cache state, which
      training never builds). The
      conv kernels already load varlen sequence starts as int64; offsets
      inside one sequence stay int32 as in 0.5.1 (upstream #1082), as for KDA
      above.
    - Backported fla PR #1109 in ops/kda/chunk_bwd.py (+ IS_NVIDIA_SM100 in
      utils/_device.py and utils/__init__.py, each marked with an inline
      "Unsloth: backported from fla PR #1109" comment): on SM100 the
      chunk_kda_bwd_kernel_wy_dqkg_fused autotune drops the BK=32 configs with
      4 or 8 warps. Reproduced on B200 with torch 2.7.1 / triton 3.3.1: every
      such config hits an illegal memory access, so the default KDA backward
      crashes while autotuning. Narrowed by Unsloth to triton < 3.4: on B200 the
      default KDA backward and the forced BK=32 / 4-warp config both run cleanly
      on triton 3.4.0, 3.5.1, 3.6.0 and 3.7.1 (torch 2.8.0 to 2.13.0), and on
      3.7.1 the autotuner picks one of those configs, so excluding them there
      cost up to 2.7% of a KDA layer fwd+bwd (B 1, T 8192, H 64, D 128).
    - Backported fla PR #1229 in ops/common/chunk_o.py (+ inline "Unsloth:
      backported from fla PR #1229" comment): on Blackwell the chunk_fwd_kernel_o
      autotune drops its 8-warp BK=BV=128 config, which hits a Triton tl.dot
      recurrence race. Reproduced on B200 with triton 3.7.1: with cu_seqlens and
      48 or 64 value heads the autotuner picks that config, and o differs from
      the dense path by up to 0.08 and between identical runs.
    - Sequence start offsets past 2**31 elements (fixed upstream after 0.5.1 by
      fla PR #1329, which forces int64 program ids everywhere). In
      ops/gated_delta_rule/chunk_fwd.py, ops/kda/{chunk_bwd,chunk_intra,
      chunk_intra_token_parallel,gate,wy_fast}.py, ops/utils/cumsum.py,
      ops/common/chunk_delta_h.py and ops/common/chunk_h.py (the simple_gla /
      GLA state kernels), kernels that add `bos * H * K`-style
      sequence start offsets in int32 take a LONG_INDEXING constexpr that
      widens `bos` to int64 only when the launcher's shapes can pass 2**31
      elements (LONG_INDEXING_THRESHOLD in utils/_config.py, exported from
      fla.utils), so normal sizes keep the int32 code (forced int64 is slower).
      Without it a packed batch or a batch whose later sequences start past
      2**31 / (heads * 128) tokens silently returned wrong outputs and
      gradients (KDA: non-finite). Offsets inside one sequence are unchanged:
      a single KDA sequence past 2**31 / (heads * 128) tokens still indexes
      in-chunk raw pointers in int32, as in 0.5.1.
    - Block pointers removed (fla PR #1062, after 0.5.1): triton main deleted
      tl.make_block_ptr / tl.advance (triton-lang/triton#10833; both raise
      NotImplementedError), so every kernel in the 19 files that used them
      (modules/{fused_norm_gate,l2norm}.py, modules/conv/triton/kernels.py,
      ops/common/{chunk_delta_h,chunk_h,chunk_o}.py, ops/cp/chunk_delta_h.py,
      ops/gated_delta_rule/{chunk_fwd,gate,wy_fast}.py, ops/gla/chunk.py,
      ops/kda/{chunk_bwd,chunk_intra,chunk_intra_token_parallel,gate,wy_fast}.py,
      ops/utils/{cumsum,pooling,solve_tril}.py) loads and stores through plain
      pointer arithmetic with explicit masks instead. #1062 was applied as a
      three-way merge (base fla 7e46dffc, theirs fla 02e27033) so every
      Unsloth change above stays in place. Plain pointers compile on every
      triton, old and new. Block pointers computed offsets in int64, so the
      loop-derived in-chunk offsets are kept int64 as upstream did afterwards:
      fla PRs #1173 / #1188 in ops/common/{chunk_delta_h,chunk_h}.py (a fresh
      i_t_int64 rather than retyping the loop variable, which torch 2.7 /
      triton 3.3 rejects) and the #1082 o_t lines in
      modules/conv/triton/kernels.py, plus the same widening of the loop-derived
      o_t / last_idx in ops/cp/chunk_delta_h.py and of o_t in the
      chunk_global_cumsum kernels of ops/utils/cumsum.py (neither fixed
      upstream; bitwise neutral at normal sizes). Without them one sequence longer than
      2**31 / (heads * head_dim) tokens hit an illegal memory access. Also
      fla PR #1330 in ops/cp/chunk_delta_h.py: merge_fwd_bwd_kernel's autotune
      key named BT, which is not a kernel argument; triton main rejects that at
      import, which made the whole vendored tree fail to inject.
    - Dropped fla/ops/gated_delta_rule/naive.py (the only einops dependency; the
      reference implementation is unused on the fast path).
    - Dropped the three heavy tilelang kernel files
      fla/ops/common/backends/tilelang/{chunk_bwd,parallel_attn_fwd,
      parallel_attn_bwd}.py. The guarded tilelang/__init__.py wrapper is kept;
      its lazy imports of those files execute only when the optional `tilelang`
      package is installed and its backend is selected (not the default path).

file_count = 73 python modules + this MANIFEST + LICENSE
closure_measured = 43 python modules for gated_delta_rule (includes naive.py);
    vendored set drops naive.py, leaving 42, plus 5 for simple_gla (its __init__,
    chunk.py, fused_recurrent.py, common/chunk_h.py, common/fused_recurrent.py),
    plus 15 for kda (11 under ops/kda, 4 under ops/gla), plus 11 for
    ShortConvolution (10 under modules/conv, modules/backends/__init__.py).

vendored_by = Unsloth (unsloth_zoo/_vendored/fla)
