v0.5.2

What's Changed

• [NSA] Support cached inference by @mutiann in #574
• [Docs] Add sparse-attention design-space note (NSA/MoBA/DSA/MSA) by @yzhangcs in #956
• [NSA] Skip empty query tiles in selective backward (bwd_dkv) by @yzhangcs in #957
• [Agents] Add correctness-gated kernel optimization loop workflow by @yzhangcs in #959
• [Docs] Restructure CONTRIBUTING; add benchmarking, testing, and style conventions by @yzhangcs in #960
• [GDN] Restrict Blackwell fwd h kernel to 2 warps by @zhiyuan1i in #953
• [NSA] Speed up parallel backward: gather+batch bwd_dkv by @yzhangcs in #961
• [NSA] Use int64 addressing to avoid int32 overflow at large T by @yzhangcs in #964
• [Attn] Prevent int32 overflow for long sequences; tune NSA large-GQA autotune by @yzhangcs in #965
• [Ops] Tile prepare_block_csr counting-sort over query blocks by @yzhangcs in #966
• [CI] Add ascend-a2-ci workflow for Atlas A2 NPU module tests by @zheliuyu in #944
• [CI] Drop auto-installed NVCC/PyTorch on H100 runners by @yzhangcs in #969
• [Agents] Tighten fla-optimization-loop guardrails from KDA-Pilot by @yzhangcs in #970
• [Fix] Scale L2Wrap logits gradient by the upstream gradient by @Sunt-ing in #963
• [Misc] Clean up feature_map type hints and add tests by @yzhangcs in #971
• [Fix] Uninitialized rotary output on left-padded batched generation by @AvivBick in #967
• [CI] Enable backend dispatch for model tests by @zhiyuan1i in #974
• [Misc] Remove unused USE_CUDA_GRAPH by @zhiyuan1i in #954
• [CI] bump GitHub Actions to latest versions by @zhiyuan1i in #976
• [Perf] Fuse GatedDeltaNet q/k/v short convolutions in the dense no-cache path by @Costa-SM in #972
• [Ops] Add naive PyTorch reference for DeepSeek Sparse Attention (DSA) by @yzhangcs in #978
• [Ops] Propagate chunk_size through non-attention chunk kernels and add coverage by @zhiyuan1i in #935
• [Ops] Expand triton-ascend backends(fla/modules, fla/ops/utils) and upgrade CANN to 9.0.0 by @zheliuyu in #973
• [Fix] remove redundant decorators, add PowGlu kernel and pin the version of torch_npu to 2.7.1.post4 for npu. by @sunyi0505 in #982
• [Fix] Restrict delta_rule prepare_wy_repr_bwd to num_warps=2 on Hopper by @zhiyuan1i in #988
• [CI] Parallelize Ascend A2 CI tests with pytest-xdist across 8 NPUs by @zheliuyu in #987
• [Fix] Prune oversized autotune configs in GLA chunk backward by @zhiyuan1i in #989
• [TileLang] Support GVA in chunk_bwd_dqkwg by @zhiyuan1i in #975
• [CI] Run op tests one file per process to avoid runner OOM by @zhiyuan1i in #991
• [Fix] Make T a runtime arg in path_attn cumprod_householder_bwd by @zhiyuan1i in #994
• [Fix] Probe TileLang availability via find_spec to avoid per-dispatch import cost by @ShareLer in #993
• [Conv] Chunk Ascend causal_conv1d launches for 65535 grid limit by @zheliuyu in #992
• [Perf] Optimize Ascend NPU activation kernels by @sunyi0505 in #995
• [Perf] Cache find_spec lookups in backend dispatch by @zhiyuan1i in #996
• [Conv] Harden triton-ascend causal_conv1d silu/add against Ascend launch limits by @zheliuyu in #997
• [GDN] Restrict Blackwell gated delta bwd autotune by @HQuanShaWu in #1000
• [Fix] fix benchamrk issues and add ub management for npu by @sunyi0505 in #1001
• [Ops] Add triton-ascend backend for l2norm and fused_norm_gate kernels by @zheliuyu in #1002
• [Fix] [Test] Remove unused GeLU/SquaredReLU NPU kernels and add Ascend UB Manager tests by @sunyi0505 in #1006
• [GDN] Add FlashQLA backend dispatch by @Erix025 in #998
• [Fix] avoid SymInt stride reads in activations under torch.compile by @zhiyuan1i in #1008
• [Fix] Expose DPLR args including cp_context/disable_recompute in chunk_rwkv7 wrapper by @sablin39 in #1004
• [Ops] Add Gluon backend for AttnRes by @yzhangcs in #1010
• [Fix] Correct Mamba-3 decay (A==+A_floor) by @netanel-haber in #1012
• [Fix] Add get_max_length to FLA cache for transformers 5.x by @zhiyuan1i in #1009
• [Test] Skip flash-attn-dependent tests when flash-attn is not installed by @zhiyuan1i in #1014
• [Perf] Optimize triton-ascend L2Norm with row tiling on Ascend NPU by @sunyi0505 in #1036
• [Fix] Store split attention decode outputs at the correct offset by @morluto in #1031
• [Fix] Avoid concurrent LSE stores in value-split attention by @morluto in #1033
• [Refactor] Centralize generation errors for unsupported cache strategies by @morluto in #1035
• [Ops] Rewrite recompute_w_u_fwd kernel for Triton Ascend backend by @OsirisDuan in #1021
• [Fix] Support value-split Wall backward with local deltas by @morluto in #1042
• [Perf] Reuse Wall autotuning across length buckets by @morluto in #1041
• [Perf] Add LayerNormGated benchmarks and BT-tile Ascend fused_norm_gate by @sunyi0505 in #1044
• [Perf] Add opt-in TileLang RWKV6 intra kernel by @markovchain-builder in #1045
• [Perf] Fuse ds in Ascend chunk_bwd_dqkwg to avoid recomputing do@v.T by @sunyi0505 in #1048
• [Perf] Fuse Ascend chunk_fwd_o inter+intra into unified 1D core-grid kernel by @OsirisDuan in #1049
• [Cleanup] Remove Ascend chunk_fwd_o kernels superseded by #1049 by @zhiyuan1i in #1050
• [Ops] Add triton-ascend backend for KDA kernels by @zheliuyu in #1047
• [CI] Add ascend-a2-benchmark-ci and fix ascend-a2-ci by @zheliuyu in #1052
• [Attn] Support for heterogeneous hybrid-attention configurations by @Michael-RDev in #1051
• [Perf] Optimize Ascend chunk_o backward with dynamic BC tiling and contiguous g loads by @sunyi0505 in #1056
• [Perf] Optimize the performance of causal conv1D on Ascend NPU by @ChengQianqian in #1024
• [Ops] Add triton-ascend backend for attn_res kernel by @zheliuyu in #1057
• [Ops] Migrate kernels off tl.make_block_ptr / tl.advance for triton main compat by @zhiyuan1i in #1062
• [Docs] Strengthen comment discipline and agent operational guidance by @zhiyuan1i in #1058
• [Perf] Optimize Ascend chunk_delta_h bwd_dhu with dynamic BK/BV tiling and host gate precomp by @sunyi0505 in #1060
• [CI] Forbid tl.make_block_ptr/advance in lint and skip H100 CI for NPU-only PRs by @zhiyuan1i in #1063
• [Perf] Optimize Ascend prepare_wy_repr_bwd with fused KV kernel, core-grid launch, and varlen tests by @sunyi0505 in #1065

New Contributors

• @Sunt-ing made their first contribution in https://gith...

v0.5.1

What's Changed

• [Typo] Fix typos in kernel comments [skip test] by @yurekami in #856
• [KDA] Support use_beta_sigmoid_in_kernel by @zhiyuan1i in #858
• [CI] Disable backend when testing models and bump tilelang by @zhiyuan1i in #861
• [Fix] GatedDeltaNet step mode crash when conv cache is None on cold Triton autotuner cache by @CiprianFlorin-Ifrim in #860
• [Cache] FLA autotune cache by @sBobHuang in #798
• Add Mamba3 by @Phoenix8215 in #855
• [README] Add Mamba-3 news entry and models table item by @yzhangcs in #865
• [Common] Support transpose_state_layout in fused_recurrent (GLA / SimpleGLA / RetNet) by @yzhangcs in #864
• [Docs] Reorganize ENVs.md into themed sections by @yzhangcs in #866
• [Conv] Fix causal_conv1d backward producing wrong gradients for non-contiguous dy by @4332001876 in #871
• [Typo] Restore comment-typo fixes regressed by #816 [skip test] by @yurekami in #873
• [Fix] Improve fused loss gradient accumulation by @yzhangcs in #874
• [Fix] Fix dk in normalized linear attention by @ayghri in #875
• [Common] Switch chunk paths to exp2 by @yzhangcs in #867
• [GDN] Fix WY backward when gating is disabled by @yzhangcs in #863
• [KDA] Fix chunk size 32 path by @yzhangcs in #876
• [Docs] Refresh README overview by @yzhangcs in #877
• [Ops] Add AttnRes operator by @yzhangcs in #878
• [Ops] Export fused AttnRes from ops [skip test] by @yzhangcs in #879
• [Models] Apply AttnRes to model definitions by @yzhangcs in #880
• [Ops] Fuse output RMSNorm into attnres kernels by @yzhangcs in #883
• [Bench] Add torch.profiler pass and run-context header to throughput bench by @yzhangcs in #886
• [Lint] Fix benchmark_training_throughput.py [skip test] by @yzhangcs in #887
• [Ops] Switch attnres to list-input + pointer-table API by @yzhangcs in #884
• [Fix] Add functools.cache for op dispatch to reduce CPU overhead by @KevinZeng08 in #881
• [Model] Add YOCO model implementation by @Shomvel in #857
• [Fix] Enable torch.compile on Python 3.10 by @yzhangcs in #894
• [Docs] Document Ruff ignore rationale by @yzhangcs in #896
• [Fix] Define FLA package import contract by @yzhangcs in #897
• [Ops] Allocate attnres fwd output in residual dtype to reduce IO by @yzhangcs in #898
• [Fix] Tolerate configs without head_dim in throughput bench by @yzhangcs in #899
• [Enhance] Extend tensor_cache to a bounded queue by @HAOCHENYE in #895
• Add Raven model integration by @AvivBick in #900
• [Raven] Remove dead params and fix o_norm affine consistency by @yzhangcs in #901
• [Raven] Fix negative RoPE offset on left-padded prefill by @yzhangcs in #902
• [Ops] Skip None args in autotune restore_value to avoid clone crash by @kasper0406 in #893
• [Ops] Rename transpose_state_layout to state_v_first and extend it to GLA / SimpleGLA chunk by @yzhangcs in #905
• [Fix] Fix packed varlen launch by @sablin39 in #904
• [Fix] Correct RWKV7 gate output chunk mask by @sablin39 in #907
• [Ops] Replace attnres int64 ptr table with padded tensor tuple by @wyc-ruiker in #912
• [Ops] Tidy attnres fused kernels: drop DTYPE, rename BL to L2 by @yzhangcs in #914
• [Perf] Read each attnres residual once in fwd and bwd by @yzhangcs in #916
• [Bench] Fix throughput-bench crash on configs without num_heads by @yzhangcs in #917
• [Fix] Large-offset Pointer Arithmetic in Blackwell GPUs by @Costa-SM in #911
• [GDN] Support use_beta_sigmoid_in_kernel & allow_neg_eigval by @yzhangcs in #919
• [Activation] Add SigmoidGLU and fix GeLUFunction backward by @yzhangcs in #921
• [Docs] Expand contributor guidelines by @yzhangcs in #923
• [Docs] Add AGENTS.md guide for AI coding agents by @yzhangcs in #922
• [Test] Add zero-gate regime to chunk_gated_delta_rule tests by @yzhangcs in #924
• [Git] Ignore Python virtual environment directories by @yzhangcs in #925
• Make varlen index helpers torch.compile-friendly by @kasper0406 in #909
• [Test] Reorganize ops/utils tests to mirror source layout by @yzhangcs in #926
• [Ops] Feat: Add triton-ascend backend support for FLA module kernels by @zheliuyu in #927
• [GDN2] Add Gated DeltaNet 2 (GDN-2) ops, layer, and tests by @shahvishal8465 in #920
• [CI] Update H100 PyTorch test matrix by @zhiyuan1i in #929
• [CI] Add agent workflow guardrails and PR title validation by @zhiyuan1i in #930
• [Attn] Add Wall attention (per-channel multiplicative-decay) op, layer, and model by @tea-more in #928
• [Refactor] Rename gated_delta_net_v2 (GDN-2) to gdn2 and align the layer with fla conventions by @yzhangcs in #931
• [Attn] Wall: add AttnRes support and align layer/model with fla conventions by @yzhangcs in #932
• [Refactor] Split fla.utils into maintainable package by @zhiyuan1i in #933
• [CI] Gate H100 Triton baseline by environment by @zhiyuan1i in #936
• [Fix] Detect consumer Blackwell (SM120) as Blackwell by @kasper0406 in #940
• [Refactor] Decouple torch/triton backend from base deps by @kashif in #937
• [Layer] Add PowGLU (Power-Gated Linear Unit) activation by @yzhangcs in #951
• [Attn] Add Parallax (parameterized local linear attention) op, layer and model by @Yifei-Zuo in #939
• [Fix] parallel_attn: keep NV==1 in backward so dq/dk reduce over full V by @kasper0406 in #941
• [Fix] FLA cache window updates and reset behavior by @Michael-RDev in #910

New Contributors

• @yurekami made their first contribution in #856
• @CiprianFlorin-Ifrim made their first contribution in #860
• @sBobHuang made their first contribution in #798
• @Phoenix8215 made their first contribution in #855
• @4332001876 made their first contribution in #871
• @ayghri made their first contribution in #875
• @HAOCHENYE made their first contribution in #895
• @AvivBick made their first contribution in #900
• @kasper0406 made their first contribution in #893
• @sablin39 made their first contribution in https://github.com/fla-org/flash-linear-a...

v0.5.0

✨ Highlights

The main new model this release is MoBA / FlashMoBA (#840, #845), which brings Moonshot's Mixture of Block Attention into the fla family.

Beyond new models, fla is also moving from a Triton-only stack to a multi-backend one: FlashKDA is added as a new backend for KDA (#852), and TileLang is introduced for GDN, KDA, and parallel attention kernels (#827, #846, #854), with more backends to come.

What's Changed

• [CP] Fix missing bos and i_h offsets in backward gk loads by @zhiyuan1i in #781
• [KDA] Clarify gate input tracking in chunk backward by @zhiyuan1i in #785
• [GDN] Fuse kkt + solve_tril kernel & unified benchmark infrastructure by @yzhangcs in #789
• [Conv] Fix int32 overflow in conv kernel pointer arithmetic for large tensors by @tmct in #783
• [GDN] Add exp2 support across chunk kernels for improved performance by @yzhangcs in #791
• Fix parameter initialization for FSDP meta device compatibility by @yzhangcs in #793
• [GDN] Fix missing mask on off-diagonal blocks in fused kkt+s… by @yzhangcs in #794
• Fix layer_norm_bwd_kernel OOB access on high-SM GPUs by @mpurland in #795
• [Misc] Upgrade minimum PyTorch requirement to 2.7.0 by @zhiyuan1i in #801
• [Conv] Fix int32 overflow in varlen conv kernel pointer arithmetic by @tmct in #803
• [GDN] Add GVA support by @zhiyuan1i in #799
• [BugFix] Fix illegal memory access in KDA backward by dropping buggy autotune configs on Hopper by @zhiyuan1i in #807
• [CE] Add logit softcapping support to fused cross entropy by @yzhangcs in #810
• [Mamba] Remove unused arguments and update to align with mamba_ssm by @PuR3Luck in #782
• [GDN] Native GVA support: remove redundant Q/K repeat and unify head naming by @yzhangcs in #812
• [KDA] Add safe_gate/lower_bound support and improve docstrings by @yzhangcs in #814
• [GDN] Add fused gate kernel with use_gate_in_kernel support by @yzhangcs in #813
• chore: add AUTHORS, unify copyright headers, and add CI workflows by @yzhangcs in #816
• [CI] Improve benchmark outputs by @yzhangcs in #817
• [CI] Fix skip-test check failing on fork PRs by @zhiyuan1i in #821
• [CP] Enable KCP for DPLR by @zhiyuan1i in #822
• [Fix] Guard checkpoint weight re-initialization in RWKV-7, Mamba, Mamba2, and LogLinearMamba2 by @puigde in #820
• fix: register default global_scratch allocator on Blackwell GPUs by @ssubbotin in #825
• [Attn] Add sliding window attention support by @yzhangcs in #824
• [Docs] Add CONTRIBUTING.md by @yzhangcs in #830
• allow neg eigvals for delta-net by @hoedt in #832
• chore: add standalone isort config by @yzhangcs in #834
• Add autotune for causal conv update by @MARD1NO in #828
• [GDN] Add TileLang backend for chunk_bwd_dqkwg kernel by @zhiyuan1i in #827
• [Refactor] Simplify TileLang backend directory structure by @yzhangcs in #835
• [CI] Post benchmark comment via workflow_run for fork-safe PRs by @yzhangcs in #841
• [KDA] Add Grouped Value Attention (GVA) support by @yzhangcs in #833
• [Attn] Add GPT-OSS-style attention sink support by @Shomvel in #831
• [Fix] respect user-provided cu_seqlens when attention_mask is present by @yzhangcs in #842
• [Fix] flatten batched qkv in varlen cu_seqlens path by @lxr-tech in #839
• [Fix] Enforce batch-size check for varlen mode across multiple ops by @zhiyuan1i in #844
• [MoBA] Integrate MOBA and FlashMOBA by @ReyJerry in #840
• [MoBA] Follow-up: fix broken import, rename layer, add modeling, tests & docs by @yzhangcs in #845
• [GDN/KDA] Fuse gate activation into fused_recurrent kernels by @yzhangcs in #848
• [TileLang] Add fwd/bwd kernel for parallel attention by @zhiyuan1i in #846
• [CI] Add issue/PR pytest command workflow by @zhiyuan1i in #843
• [GDN] Optimize b_dg computation in chunk_bwd_kernel_dqkwg#USE_G by @MzeroMiko in #823
• [GDN][Tilelang] Optimize b_dg computation in chunk_bwd_kernel_dqkwg by @zhiyuan1i in #849
• [Linear Attention] Update fused_recurrent.py for inference with normalization by @yiyousong in #268
• [Fix] Fix incorrect cumsum dim for naive_chunk_linear_attn normalize by @zhiyuan1i in #851
• [Cleanup] Remove deprecated head_first parameter from public ops by @yzhangcs in #853
• [KDA] Support FLASHKDA backend by @zhiyuan1i in #852
• [KDA][TileLang] Add TileLang backend for chunk_kda_bwd_wy_dqkg_fused by @zhiyuan1i in #854

New Contributors

• @tmct made their first contribution in #783
• @mpurland made their first contribution in #795
• @PuR3Luck made their first contribution in #782
• @puigde made their first contribution in #820
• @ssubbotin made their first contribution in #825
• @hoedt made their first contribution in #832
• @MARD1NO made their first contribution in #828
• @Shomvel made their first contribution in #831
• @lxr-tech made their first contribution in #839
• @MzeroMiko made their first contribution in #823
• @yiyousong made their first contribution in #268

Full Changelog: v0.4.2...v0.5.0

v0.4.2

What's Changed

• [Misc] Use autopep8 to keep style by @zhiyuan1i in #697
• [Misc] Reduce D2H/H2D Sync by @zhiyuan1i in #698
• [Conv]: Support mix mode(Triton fwd and CUDA bwd) by @zhiyuan1i in #699
• [Test] Add memory guard fixtures for CUDA memory safety testing by @zhiyuan1i in #700
• [KDA] Add lowerbound gate function by @zhiyuan1i in #701
• [KDA] Remove deprecated head_first in kda gate func by @zhiyuan1i in #702
• [KDA] Speed up chunk_kda by introducing lowerbound gate by @zhiyuan1i in #703
• [NSA] fix varlen related logic in cmp dkv kernel by @yibozhong in #707
• [DPLR] Speed up DPLR by lowerbound gate by @zhiyuan1i in #709
• [Misc] Enhance non-cuda platform ci by @zhiyuan1i in #708
• [Conv] Add non-contiguous tensor support for convolution ops by @zhiyuan1i in #712
• [Conv] Fix corner case by @zhiyuan1i in #714
• [Conv] Refactor dh0 into separate Triton kernel and add gradient tests by @zhiyuan1i in #717
• [Misc] Wrap exp/log math ops with @triton.jit and enforce float32 pre… by @zhiyuan1i in #720
• [Conv] Refractor to packages by @zhiyuan1i in #722
• [Conv] Clean duplicate calculate chunk_indices by @zhiyuan1i in #724
• [DPLR] Add disable_recompute support for DPLR chunk op by @zhiyuan1i in #726
• [CP] fuse fwd/bwd kernels and fix IMA in long context by @zhiyuan1i in #733
• [KCP] add KCP.md; fix fp32 precision in M matrix chain; cleanup CP tests by @zhiyuan1i in #740
• [Backend] Introduce dispatch system by @zhiyuan1i in #741
• [Backend] Select from available backends based on priority order by @zhiyuan1i in #742
• [MAMBA2] fix initialization for mamba2 by @mayank31398 in #739
• [KDA] Refractor interface by @zhiyuan1i in #744
• [Deltarule] Added intra-card context parallel optimization for KDA and GDN by @zhiyuan1i in #743
• [Misc] centralize reference implementations in naive.py by @zhiyuan1i in #746
• [Cache] Fix get_seq_length to return per-layer length by @zhiyuan1i in #748
• fewer l2norm recompilations by @tyler-romero in #745
• [OJA] Integrate Gated OJA Rule by @AwesomeSeq in #730
• [Misc] remove redundant dot precision param in KDA recompute_w_u by @KevinZeng08 in #750
• Fix shared memory guards for AMD RDNA GPUs (64KB shared mem) by @Gildoniel in #751
• [Fix] Guard A_log and dt_bias re-initialization against loaded checkpoint values in GatedDeltaNet, Comba, and KDA by @ljxw88 in #754
• [MAMBA-2] fix mamba-2 init for FSDP-2 with DTensors by @mayank31398 in #753
• [Deltrarule] Add cache for intra cp by @zhiyuan1i in #755
• Update Windows warning to detect triton-windows by @erm14254 in #757
• [Misc] Reduce recompile by @zhiyuan1i in #764
• [PaTH] Prevent int32 index overflow for long sequences by @zhixuan-lin in #769
• [Test] Add regression tests for cache seen-token bug (GH-766) by @zhiyuan1i in #768
• [Misc] removes all @torch.jit.script decorators from the codebase by @zhiyuan1i in #767
• [Conv] Fix invalid memory access by @zhiyuan1i in #774
• [KDA][GDN] Support transpose_state_layout for [V,K] state memory layout by @zhiyuan1i in #776
• [GDN] Enhance Triton 3.2 compatibility by @winglet0996 in #773
• [Model] Unify cache function by @zhiyuan1i in #777

New Contributors

• @mayank31398 made their first contribution in #739
• @tyler-romero made their first contribution in #745
• @KevinZeng08 made their first contribution in #750
• @Gildoniel made their first contribution in #751
• @ljxw88 made their first contribution in #754
• @erm14254 made their first contribution in #757
• @winglet0996 made their first contribution in #773

Full Changelog: v0.4.1...v0.4.2

🎄 v0.4.1

What's Changed

• [GDN] fix oom on A6000 by @sustcsonglin in #622
• [KDA] Fix _no_weight_decay term by @yzhangcs in #629
• Lint all files by @zhiyuan1i in #628
• Fix kda.gate not enforcing contiguous mem layout by @Qubitium in #627
• Fix docstring for 'g' parameter shape by @GuoYiFantastic in #631
• Fix: Correct K/V dimension mismatch in path_attn bwd kernels"changing K/BK to V/BV for v and dv operations by @ReyJerry in #633
• KDA - fix: don't force fused_recurrent when in training mode with small sequences by @masc-it in #636
• [KDA]: Fuse beta.float().sigmoid() in fused_kda_gate by @zhiyuan1i in #642
• [Perf] add chunk_indices parameter to avoid redundant computation by @zhiyuan1i in #641
• Added a badge for Ask DeepWiki to the README to auto-refresh the wiki weekly by @richardodliu in #644
• [Tril] Enable precision autotune[skip test] by @zhiyuan1i in #646
• [BC] Capitalize all envs by @yzhangcs in #650
• deprecate fused_chunk_gla and safe_exp; fix kda exp mask by @sustcsonglin in #652
• [kda kernel optimization] implement token-parallel intra-chunk attention by @sustcsonglin in #653
• [KDA] Faster inter computation in 64x64 intra fwd by @yzhangcs in #658
• Add PTX softplus by @yzhangcs in #660
• [KDA] Support fused forget gate by @yzhangcs in #662
• [KDA] Remove beta from fused gate by @yzhangcs in #665
• [Softplus] Support AMD/Intel devices by @zhiyuan1i in #664
• [KDA] Fix ood bugs in intra fwd by @yzhangcs in #673
• [L2Norm] Avoid recompilation for variable-length inputs by @retonym in #669
• fix: set the dtype of RMSNorm to float32 to avoid precision underflow by @pprp in #676
• [KDA] Changed all exp to exp2 by @Nathancgy in #679
• [KDA] Fuse inplace add by @yzhangcs in #682
• Add head_dim parameter to NSA layer by @mutiann in #683
• [KDA] Fuse dAqk and dv by @yzhangcs in #689
• Temporary workaround to disable TritonGPUHoistTMEMAlloc in b_dk += tl.dot(tl.trans(b_dA), b_kb) by @rucnyz in #687
• [NSA] fix compression branch dkv kernel dk and dv pointer impl by @yibozhong in #690
• [KDA] fused bwd kernels inter and prepare wy by @Nathancgy in #688
• [GDN] Fix potential ood for long inputs by @yzhangcs in #692
• [GDN] Support beta in float32 by @AwesomeSeq in #693
• [GSA] Fix gate oob bugs by @yzhangcs in #694

New Contributors

• @Qubitium made their first contribution in #627
• @GuoYiFantastic made their first contribution in #631
• @ReyJerry made their first contribution in #633
• @masc-it made their first contribution in #636
• @retonym made their first contribution in #669
• @pprp made their first contribution in #676
• @mutiann made their first contribution in #683

Full Changelog: v0.4.0...v0.4.1

v0.4.0

🧠 New Models

• 🌑 Kimi Delta Attention by @yzhangcs @sustcsonglin

What's Changed

• [GDN] Fix tiling bugs once gv applied by @yzhangcs in #589
• [Conv] Add comprehensive docstring and change default backend to triton by @zhiyuan1i in #592
• Update cumprod_householder_bwd.py by @SeepingFragranceLock in #593
• [FIX] Correct cumsum dimension in normalize_output by @sirluk in #594
• [Triton] Add autotune caching support for Triton kernels by @zhiyuan1i in #598
• [DeltaFormer] Add Model by @Nathancgy in #585
• [DeltaFormer] Replace GenerationMixin with FLAGenerationMixin and upd… by @zhiyuan1i in #600
• [Cache] Fix from_legacy_cache by @zhiyuan1i in #605
• [Deps] Make pytest an optional dependency by @wedaly in #610
• [DeltaFormer] Fixed testing ops error by @Nathancgy in #602
• [Conv] Fix potential OOB problems by @yzhangcs in #615
• [Deps] Minimize deps by @zhiyuan1i in #617
• Determine the chunk size at the kernel entry by @yzhangcs in #619
• Add KDA by @yzhangcs in #621
• [Lint] Migrate from flake8/isort to ruff for faster linting by @zhiyuan1i in #613

New Contributors

• @SeepingFragranceLock made their first contribution in #593
• @sirluk made their first contribution in #594
• @Nathancgy made their first contribution in #585
• @wedaly made their first contribution in #610

Full Changelog: v0.3.2...v0.4.0

v0.3.2

📣 Highlights

Starting with this release, every time we ship a new version of flash-linear-attention, we will simultaneously publish fla-core: a minimal-dependency subset of the main repo that contains only the essentials.

🧠 New Models

• 🌲 Log Linear Attention by @2022tgoel

What's Changed

• [Conv] Provide fn interface for causal_conv1d by @yzhangcs in #578
• [Log Linear Attention] add backward pass by @2022tgoel in #577
• [PaTH] Fix q init & dq masking by @yzhangcs in #581
• [TokenShift] Fix a bug in decoding by @zhiyuan1i in #583
• [Deps] Lock transformers<4.56.0 by @zhiyuan1i in #582
• [Log-Linear Attention] add models by @2022tgoel in #579
• [Deps] Upgrade to transformers 4.56.x by @zhiyuan1i in #587
• [Build] Split package distribution[skip test] by @zhiyuan1i in #588

Full Changelog: v0.3.1...v0.3.2

v0.3.1

What's Changed

• [Misc] Change grid to support long ctx by @zhiyuan1i in #528
• [RWKV7] Reduce CPU overhead by @zhiyuan1i in #529
• [Tokenshift] Support SP and cache by @zhiyuan1i in #531
• [RWKV7] Use tokenshift to save cache by @zhiyuan1i in #532
• [RWKV7] Fix the issue of RWKV7 initialization with BFloat16 data type on CPU. by @zhiyuan1i in #538
• [CI] Add compatibility check by @zhiyuan1i in #536
• [ShortConv] Support cache in prefill by @zhiyuan1i in #535
• [WIP] Add Log-Linear Attention by @2022tgoel in #524
• [Cache] Upgrade to transformer>= v4.48[skip test] by @zhiyuan1i in #541
• [Misc.] Set env var TRITON_F32_DEFAULT to ieee when tf32 is not supported on NVIDIA by @KevlarKanou in #544
• [CI] Fix mirror for building triton by @zhiyuan1i in #543
• Log-Linear Attention Tests by @2022tgoel in #542
• [CI] Add proxy config for git by @zhiyuan1i in #548
• [Conv] Fix warning issue by @zhiyuan1i in #549
• [Misc.] Eliminate recompilation in layer-norm kernels caused by dynam… by @zhiyuan1i in #545
• [Misc.] Add activations for non-cuda Backends by @zhiyuan1i in #174
• [TMA] Accelerate solve_tril with TMA descriptors[skip test] by @zhiyuan1i in #550
• [CI] Upgrade to latest casual-conv1d and fix triton build for 3.4.x by @zhiyuan1i in #551
• [CI] Fix support for Intel GPU by @zhiyuan1i in #554
• [Fix] Fix Triton Error for HeadDim < 16[skip test] by @zhiyuan1i in #556
• [GLA] Fix simple_gla Test by @zhiyuan1i in #558
• [CI] Fix CI script errors[skip test] by @zhiyuan1i in #566
• require transformers <= 4.53.3 by @richardodliu in #570
• [Deps] Adopt transformers>4.53.3 by @zhiyuan1i in #571
• [Misc.] Clean codes and make mypy happy by @zhiyuan1i in #572
• [Models]: Add MoM by @WKX933 in #442
• [MoM]Fix lint by @JusenD in #573
• [Refactor] Apply GradientCheckpointingLayer to all model layers by @yzhangcs in #575
• [Mamba] Fix errors in Triton backend by @zhiyuan1i in #576

New Contributors

• @2022tgoel made their first contribution in #524
• @KevlarKanou made their first contribution in #544
• @richardodliu made their first contribution in #570
• @WKX933 made their first contribution in #442

Full Changelog: v0.3.0...v0.3.1

v0.3.0

Highlights

🧠 New Models

We are excited to expand our model library with the addition of four powerful new architectures.

• 🎉 MesaNet by @sustcsonglin
• 🛣️ PaTH by @sustcsonglin
• 🐍 Comba by @AwesomeSeq @yzhangcs
• 🐳 MLA by @toothacher17 @yzhangcs

What's Changed

• [MesaNet] add kernel impl. by @sustcsonglin in #419
• [GDN] Add support for inference with GVA by @yzhangcs in #429
• [HGRN] remove unused q_conv1d by @yibozhong in #430
• Update mesa_net.py by @jovoswald in #434
• [Gated DeltaNet] Refactor the kernel to remove one matrix inversion by @sustcsonglin in #433
• [Modules] Add L2Warp to maintain bf16 precision by @zhiyuan1i in #438
• [RWKV]: Set default scale to None by @zhiyuan1i in #445
• [Typos] Change scale docs to (Optional[float]) [skip test] by @zhiyuan1i in #446
• [Modules] Enhance Testing of l2warp by @zhiyuan1i in #448
• [CI] Upgrade CI envs to torch~=2.7.0 by @zhiyuan1i in #450
• [Mesa] misc. fix by @sustcsonglin in #449
• [Models]: Add Comba Implementation by @AwesomeSeq in #444
• [Test] Walk around the bug of causal_conv1d by @zhiyuan1i in #453
• [Utils] Add deprecation handling for kwargs with deprecate_kwarg decorator by @yzhangcs in #455
• [ShortConv] Replace use_fast_conv1d with backend parameter by @yzhangcs in #456
• [Docs] Update tensor shape descriptions and deprecate head_first argument by @yzhangcs in #457
• [Simple GLA] Support dg when dht passed by @yzhangcs in #459
• [Mesa] Improve precision by @sustcsonglin in #460
• [Comba] Remove problematic safe_exp by @yzhangcs in #466
• [TokenShift] Fix invalid argument on AMD GPUs by @zhiyuan1i in #464
• [Test] Refractor model testing[skip test] by @zhiyuan1i in #467
• [Testing] Enhance generation testing by @sustcsonglin in #468
• [Simple GLA] Remove unnecessary dg for data-independent decay by @yzhangcs in #469
• [CI] Update workflow by @zhiyuan1i in #473
• [Misc.] Enhance support for some platforms by @zhiyuan1i in #470
• [Gated Delta Product] Optimize kernels by @sustcsonglin in #472
• [README] Add support for aarch64 by @zhiyuan1i in #475
• [Cache] Fix bad seen_tokens update by @yzhangcs in #478
• [CI] Revert causal-conv1d to 2a288a1 by @zhiyuan1i in #480
• [Parallel] Fix all tokens offsets by @yzhangcs in #479
• Use tl.exp2 for all gating operations by @yzhangcs in #361
• Refactor modeling tests by @yzhangcs in #482
• Add L2_norm for p in Recurrent ops to fix generation error by @AwesomeSeq in #483
• Refactor benchmark: adapt to latest FLA benchmark interface by @yuweih205 in #488
• [GLA] Remove all safe_exp ops by @yzhangcs in #489
• [MesaNet] Remove all safe_exp ops & Refactor tests by @yzhangcs in #490
• [Misc.] Support PT2.5 by @zhiyuan1i in #491
• [Misc.] Fast testing & Autotune by @sustcsonglin in #476
• fix: update import path for causal_conv1d by @yuweih205 in #492
• Make RWKV-7 init match official RWKV-LM by @johanwind in #493
• Modernize the fused_chunk impls by @yzhangcs in #437
• [ShortConv] Fix bad conv weight input shape during inference by @yzhangcs in #495
• [DeltaProduct] chore: remove unused functions by @timurcarstensen in #496
• [CI] Fix pipeline in GPU CIs by @zhiyuan1i in #497
• [RWKV] Make torch.compile decorator compatible with python3.10 by @zhiyuan1i in #498
• [GDN] Fuse 64x64 matrix inverse kernel by @yzhangcs in #501
• [L2Norm] Speedup by saving rstd by @yzhangcs in #506
• [Norm] Move eps out of sqrt by @yzhangcs in #508
• Correct types of constructor arguments with issues for configuration classes by @V0XNIHILI in #509
• Fix typo: suppoerted -> supported by @zxytim in #510
• [RWKV7] Increase Lora shape for headdim>64 by @zhiyuan1i in #512
• [Delta Rule] Support gk for WY reprs by @yzhangcs in #514
• [PaTH attention] Support headdim 128 & refactor kernel for better stability by @sustcsonglin in #503
• [Rotary] Fix max_seqlen under varlen mode by @yzhangcs in #516
• [Misc] Skip testing models on Nvidia 4090 CI by @zhiyuan1i in #517
• [GDP] Delete duplicated code by @yzhangcs in #518
• [WIP] Add MLA layers into fla by @toothacher17 in #395
• [Mamba] Add triton conv1d backend and fix mamba2 test by @zhiyuan1i in #520
• [Typo] Fix types in all configuration files[skip test] by @V0XNIHILI in #513
• [GSA] Fix memory boundary conditions by @JusenD in #527

New Contributors

• @jovoswald made their first contribution in #434
• @AwesomeSeq made their first contribution in #444
• @yuweih205 made their first contribution in #488
• @V0XNIHILI made their first contribution in #509
• @zxytim made their first contribution in #510
• @toothacher17 made their first contribution in #395
• @JusenD made their first contribution in #527

Full Changelog: v0.2.2...v0.3.0

v0.2.2

What's Changed

• [TokenShift] support fused_token_shift with varlen by @zhiyuan1i in #373
• [Mamba] Use official init strategies by @yzhangcs in #374
• [Mamba2] Create attn layer by @yzhangcs in #375
• [Mamba] Add attn layer & fix configs by @yzhangcs in #376
• [RWKV7] Update fused_addcmul impls by @zhiyuan1i in #378
• [RWKV7]: Rewrite docs to match Triton codes. by @zhiyuan1i in #381
• [RWKV7] Fix convert script by @zhiyuan1i in #383
• [Misc.] Update triton-nightly.yml by @zhiyuan1i in #382
• [PaTH] Add PaTH attention model and kernel by @sustcsonglin in #384
• [Tests] Enable tests with causal_conv1d on H100 CIs by @zhiyuan1i in #385
• [GDN]: initializing A_log and dt_bias in _init_weights by @HanGuo97 in #380
• [Utils] Add fused pack/unpack fns by @yzhangcs in #386
• [RWKV7] Strictly initialize rwkv7 according to RWKV-LM by @zhiyuan1i in #387
• [chore] switched to processing_class kwarg inside Trainer invocation by @timurcarstensen in #391
• [RWKV7] Update initialization to sync with latest RWKV-LM by @zhiyuan1i in #393
• [Token Shift]: Fix potential cuda kernel parameter error for varlen by @zhiyuan1i in #397
• [DeltaProduct] fix query conv cache, remove extraneous query convs by @timurcarstensen in #396
• [Misc.] Log warnings when Triton is older than 3.2.0 by @zhiyuan1i in #394
• [RWKV7]: clean fused_addcmul_rwkv7 impls by @zhiyuan1i in #404
• [README] Update FoX venue info by @zhixuan-lin in #406
• Added details to some formulas, fixed the display error of the L2 Loss formula by @Beortext in #407
• [RWKV7] Change fp32 errors to warnings by @zhiyuan1i in #412
• [Misc.] Add exist_ok=True to all models by @zhiyuan1i in #413
• Add Rodimus impl into fla by @ziHoHe in #416
• Align RWKV7 LoRA Rank Initialization with official Implementation by @WuTianyi321 in #418
• [Canon] Add triton impls by @yzhangcs in #388
• [GDN] Support Gated Value Attention (GVA) by @Rafa-zy in #421
• [RWKV7]: clean some imps by @zhiyuan1i in #420
• [RoPE] Fix out-of-boundary bugs by @yzhangcs in #423
• [RWKV] Fix cu_seqlens with gradient checkpoint by @zhiyuan1i in #422

New Contributors

• @timurcarstensen made their first contribution in #391
• @ziHoHe made their first contribution in #416
• @WuTianyi321 made their first contribution in #418
• @Rafa-zy made their first contribution in #421

Full Changelog: v0.2.1...v0.2.2