v0.2.0.post2

What's Changed

• ci: fix the update_whl_index script to regonize version number with "post" and add torch2.5 by @yzh119 in #694
• bugfix: casting int array to int32 for rope input arguments by @yzh119 in #697
• bugfix: only use sm90 group gemm when torch cuda >= 12.3 by @yzh119 in #699
• misc: remove release-please workflow by @yzh119 in #705
• Customizable SM90 prefill kernels. by @hyhieu in #704
• hotfix: revert torch.library register by @yzh119 in #709
• Improve compatibility with pytorch 2.5 by @zifeitong in #711
• misc: add bibtex reference by @yzh119 in #712
• sampling: simplify min-p sampling by @yzh119 in #713
• perf: fix the iteration bound of SWA in FA2 prefill template by @yzh119 in #714
• bugfix: fix min-p AOT compilation in #713 by @yzh119 in #717
• Triton implementation of silu_and_mul by @nandor in #716
• bugfix: FusedAddRMSNorm kernels might require more than 48KB shared memory when d is large. by @bobboli in #718
• bugfix: Choose sm90 kernels only for Hopper GPUs. by @bobboli in #719
• Finer-grained control over fp16/fp8 builds by @nandor in #722
• Align KV chunk size binary search with actual KV chunk splitting. by @timzsu in #728
• ci: rename python package name to flashinfer-python by @yzh119 in #729
• Add a note about int32/int64 datatypes to the kv_layout tutorial by @fergusfinn in #737
• fix return type of cuBLAS by @zhyncs in #749
• [Refactor] Unify JIT/Customization/AOT mode by @yzh119 in #748
• Move allocations out of torch ops by @nandor in #740
• [Lint] Fix some linting issues and provide automatic format check script by @LeiWang1999 in #743
• Filter out unsupported head dim for sm90 by @abcdabcd987 in #751
• bugfix: various AOT issues by @abcdabcd987 in #752
• [bugfix] Fix cpp tests/benchmarks by @yzh119 in #753
• fix pin memory device by @youkaichao in #755
• Add dev container for easier development by @ByronHsu in #680
• hotfix: bugfix to #756 by @yzh119 in #757
• Change apply_rope_with_cos_sin_cache to accept cos_sin_cache by @ByronHsu in #754
• fix: match statement not supported in Python 3.8 by @xslingcn in #759
• bugfix: use actual sm count for num_sm90_ctas by @LLLLKKKK in #762
• bugfix: Fix block-sparse attention API by @yzh119 in #767
• Version bump: v0.2.0.post2 by @yzh119 in #768

New Contributors

• @hyhieu made their first contribution in #704
• @zifeitong made their first contribution in #711
• @bobboli made their first contribution in #718
• @timzsu made their first contribution in #728
• @fergusfinn made their first contribution in #737
• @LeiWang1999 made their first contribution in #743
• @youkaichao made their first contribution in #755
• @LLLLKKKK made their first contribution in #762

Full Changelog: v0.2.0.post1...v0.2.0.post2

v0.2.0.post1

0.2.0.post1 (2024-12-22)

Bug Fixes

• bug fix on determine_attention_backend condition (#688) (bcf7a3e)
• accelerate plan speed of fa3 template (#690) (db8f04d)

v0.2.0

0.2.0 (2024-12-17)

Release Blog.

Features

• add rotary_dim argument to rope APIs for partial apply rope (#599) (eb9bc71)
• add a use_softmax field in variant class (#533) (d81af97)
• add an option non_blocking to plan function (#622) (560af6f)
• add gemma_rmsnorm and gemma_fused_add_rmsnorm (#477) (1a6b17e)
• add group size 3 to GQA decode dispatch (#558) (6227562)
• add JIT compilation support for FA3 templates (#672) (d4e8d79)
• allow the cascade kernels to be executed using varying sequence lenghts (#627) (92ac440)
• CUDAGraph compatibility of multi-level cascade inference APIs (#586) (2332e8a)
• fix the maximal grid dimension in prefill planning with CUDA graphs (#639) (86ca89a)
• improve the precision of the FusedAddRMSNormKernel function (#587) (c7dc921)
• JIT compilation (#507) (3613a5b)
• modify group-gemm stage number (#497) (52dab1d)
• non-contiguous query with paged kv cache (#553) (89f2c4a)
• pass a dynamic token count to the cascade kernels (#635) (5fe9f7d)
• simplify prefill JIT compilation (#605) (fe4f898)
• specify gemm backend (#648) (0cc1a51)
• support cached cos/sin in rope APIs (#585) (83e541d)
• support huggingface transformer style rope interface (#568) (4f40420)
• support sm90 cutlass group gemm (#509) (794bdda)
• torch custom_op fix for rope (#569) (3e104bc)
• torch custom_op support: norm (#552) (f6e0010)
• torch.compile and custom_op support (#554) (9bf916f)
• warmup for jit kernel tests (#629) (8f5f349)

Bug Fixes

• AOT compiler flags on non-sm90 (#522) (0aa4726)
• batch decode kernel redundant store output to gmem (#505) (90e42a7)
• compatible with torch 2.2 (#478) (ac41d1b)
• #452 (b53a46f)
• remove redundant load (#495) (2de16b0)
• update bmm fp8 test (#487) (45eac04)

Performance Improvements

• accelerate JIT compilation speed (#618) (eaf73fd)
• Dense and sparse customizable flashattention-3 template (#667) (51236c9)
• fix prefill kernel performance degradation (step 1) (#602) (595cf60)
• fix the performance issue of append_paged_kv_cache (#588) (e15f7c9)
• improve parallelism in RoPE with pos_ids (#609) (ff05155)
• improve plan performance by using non-blocking memcpy (#547) (41ebe6d)
• reduce the read and write of shared memory in the FusedAddRMSNormKernel (#592) (2043ca2)
• reduce total_num_tiles_q by one (#644) (553ace5)
• remove unnecessary contiguous operation in block sparse attention (#561) (7a7ad46)
• speedup jit compilation of prefill attention kernels (#632) (a059586)
• use cuda-core implemention for io-bound block-sparse attention (#560) (3fbf028)

v0.1.6

0.1.6 (2024-08-27)

SM75 Support

Starting from 0.1.6, our pre-built wheels include experimental support sm75 (Turing architecture GPUs such as Tesla T4, Quadro RTX 6000 and RTX 2080).

API Changes

plan/run

Since 0.1.6 on, begin_forward/forward/end_forward APIs are replaced with the new plan/run API.

• forward is renamed to run, which is more precise and consistent with the naming convention of cutlass's python API.
• begin_forward is renamed to plan, which is consistent with the naming convention of nvmath API.
• end_forward is deprecated and has no effect after this PR.

There is some slight difference between the old forward and the new run API:

• All extra arguments such as causal and logits_soft_cap will be provided in plan (previously begin_forward) API, and cached until next plan call, and we only need to provide query and KV-Cache tensors in run API.

The old begin_forward/forward/end_forward APIs are still functional, but we will gradually deprecate them in future releases.

Check #466 for more details.

MultiLevelCascadeAttentionWrapper

Since 0.1.6 on, we introduce a new MultiLevelCascadeAttentionWrapper API for cascade inference,
which supports multi-level cascade inference where all levels' KV-Cache can be managed in a unified Paged KV-Cache.

See documentation and tutorial on API usage and layout explaination.

The old BatchDecodeWithSharedPrefixPagedKVCacheWrapper and BatchPrefillWithSharedPrefixPagedKVCacheWrapper will be deprecated in future releases.

Features

• sm75 support (#448, #449)
• add MultiLevelCascadeAttentionWrapper API (#462) (1e37989)
• add accept num, emit num metric for ChainSpeculativeSampling (#450) (fa38b5e)
• support bmm fp8 (#469) (f1c0b68)

Refactor

• refactor: replace begin_forward/forward/end_forward with plan/run #466

Misc

• misc: improve error handling of sampling kernels (#456) (0dce178)

Performance Improvements

• slight optimization on f16->f8 fragment layout swizzling (#453) (0d61871)
• slight optimization on fragment layout swizzle (#458) (7c397cb)
• use persistent kernel for merging attention states (#459) (be6bf5b)

Acknowledgement

We thank @LiuXiaoxuanPKU on enhance of speculative sampling operator, @merrymercy on API change suggestion and @zhyncs on integrating fp8 BMM cublas implementation.

v0.1.5

0.1.5 (2024-08-13)

Bugfix

• Fix PagedPrefill python api and some typos (#441) (3fff008)
• fix prefill kernels' lse result for empty kv-cache (#440) (6ac28f4)

Features

• decouple float and int workspace buffer (#442) (a7ee566)

Performance Improvements

• faster fp8->fp16 dequantization for pre sm_90 arch (#439) (c93f647)

Acknowledgement

We thank contributions and feedbacks from the community: @comaniac, @hnyls2002, @jianfei-wangg, @Yard1.

v0.1.4

0.1.4 (2024-08-09)

Features

• append attention kernels for fp8 kv-cache (#420) (906c2f5)
• support min_p sampling (#422) (d52f2da)
• deterministic sampling (#417) (0dd801d)
• more sampling operator options (#431) (68df9c4)
• support fused add rmsnorm (#419) (b781513)
• support fused silu mul (#427) (ea0ba9a)
• feat: support fused gelu tanh mul (#434) (2c9d1c3)

Bug Fixes

• fix dispatch fp16 type when enable fp8 (#430) (daa5566)
• improve numerical stability of sampling kernels (#429) (898d8ea)

Other improvements

• break up _kernels into multiple modules (#428) (8e482d9)

Acknowledgement

We thank contributions and feedbacks from the community: @comaniac, @esmeetu, @LiuXiaoxuanPKU, @peng1999, @xslingcn, @Yard1, @zhyncs.

v0.1.3

0.1.3 (2024-07-31)

Bugfix

• bugfix: Fix cudagraph mode of BatchPrefillWithRaggedKVCacheWrapper (#412) (9907bc)
• fix cu118 cub usage for sampling kernels (#410) (58d359)

Misc

• enhance allocator error info and add shape check for prefill begin forward functions (#413) (5e36c5)

v0.1.2

0.1.2 (2024-07-29)

Bugfix

• Fix the sampling kernel bug for cu118 (#386, #387) (0cd499, dc3f18)

Features

• add llama 3.1 style rope (#401) (4c89dec)
• non-inplace rope operators (#405) (74ffba1)
• sliding window attention (#406) (28cffd3)
• support non-contiguous (packed) input for prefill kernels (#404) (68c3719)

Performance Improvements

• slight optimization on merge states (#313) (701c813)

v0.1.1

0.1.1 (2024-07-20)

Bugfix

• fix the invalid kernel configuration for architectures with small shared memory size (#385) (cdac57)

Features

• expose decoupled kv-cache to pytorch api (#383) (457a0ae)

Performance Improvements

• use stmatrix in epilogue for sm90+ (#380) (c6f20d1)

v0.1.0

0.1.0 (2024-07-17)

Features

• Add mask to merge_state_in_place (#372) (e14fa81)
• expose pytorch api for block sparse attention (#375) (4bba6fa)
• Fused GPU sampling kernel for joint top-k & top-p sampling (#374) (6e028eb)