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 and the GLM-5.3-Flash / Kimi-Linear Kimi delta
    attention models without requiring `pip install flash-linear-attention`.

exported_api =
    fla.modules.FusedRMSNormGated
    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.
    - 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).
    - 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 = 62 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).

vendored_by = Unsloth (unsloth_zoo/_vendored/fla)
