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)

v0.0.9

0.0.9 (2024-07-12)

Bugfix

• fix the decode kernel segfault in cudagraph mode (#368)(c69cfa)

• fix decode kernels output for empty kv cache (#363)(ac72b1)
• check gpu id in PyTorch APIs and use input tensor's gpu default stream (#361)(1b84fa)

Performance Improvements

• accelerate alibi (#365) (4f0a9f9)
• accelerate gqa performance (#356) (e56ddad)
• Optimize tensor conversions in C++ code to avoid unnecessary copies (#366) (1116237)

Acknowledgement

We thank @Yard1, @Ying1123 and @zhyncs for their contributions.

v0.0.8

0.0.8 (2024-07-03)

Bugfix

• fix prefill/append kernel behavior for empty kv-cache (#353) (7adc8c)
• fix decode attention kernel with logits cap (#350) (f5f7a2)

v0.0.7

0.0.7 (2024-06-28)

Breaking Changes

• batch_decode_with_padded_kv_cache was removed, we encourage user to use BatchDecodeWithPagedKVCacheWrapper instead. (#343)

Bugfix

• fix the forward_return_lse function in BatchPrefillWithRaggedKVCache class (#337)
• fix the scheduler behavior of large page size (#333)

Features

• customize logits_soft_cap value (#339) (a2498f5)

Performance Improvements

• change minimal kv_chunk_size back to 128 (#329) (f237f5f)
• more options for kv tile size (#336) (bf2a6c7)

v0.0.6

0.0.6 (2024-06-21)

Performance Improvements

• use 1x4 warp layout for small query length (not activated because of large binary size) (#322) (4e89b4d)

v0.0.5

0.0.5 (2024-06-20)

Highlights

• Support any GQA group size for tensor-cores kernels.
• Support any page size for tensor-cores kernels.
• Support CUDA-Graph for prefill/decode APIs.
• Add an option to accelerate decode kernels with Tensor Cores.
• Support custom attention mask. (https://docs.flashinfer.ai/tutorials/kv_layout.html#mask-layout-2d-ragged-tensor)
• Support logits cap in Grok-1 models.
• Fused GPU-sampling kernels: top-p, top-k, speculative verification. (https://docs.flashinfer.ai/api/python/sampling.html)
• PyTorch wrapper of group-gemm cutlass kernels. (https://docs.flashinfer.ai/api/python/group_gemm.html)

Acknowledgement

We thank @ibsidorenko, @LiuXiaoxuanPKU, @Yard1 @AgrawalAmey, @xuzhenqi, @mgerstgrasser, @esmeetu, @yz-tang, @HSQ79815, @Qubitium, @shreygupta2809, @sighingnow, @vinx13, @tqchen, @merrymercy, @comaniac and many others for their contributions and helpful discussions for 0.0.5 release.

Refactor

• support any GQA group size for tensor-cores kernels (#301) (c111ca)
• support any page size for tensor-cores kernels (#306) (82fd8c)

Features

• add use_tensor_cores option to decode kernels to accelerate GQA (#317) (3b50dd5)
• add group gemm operators (#282) (e08ba42)
• initial support of distributed operators (#289) (03553da)
• initial support of logits hook (#298) (ab1e2ad)
• Separate Q and KV dtypes for decode (#286) (5602659)
• support cuda graph for batched multi-query(prefill/append) attention (#275) (83ceb67)
• support cuda graph for batched multi-query(prefill/append) attention (#277) (24cc583)
• support custom attention mask in prefill/append attention kernels (#266) (7304282)
• fused speculative sampilng kernels (#259) (cea2bb)
• expose sampling APIs in pytorch (#238) (092902)

Performance Improvements

• initial cuda graph support (#256) (7e9cc7f)
• split kv-cache for prefill/append kernels (#310) (f0bb0a3)
• use packed bit array for attention mask (#308) (3d43dc9)

v0.0.4

0.0.4 (2024-05-01)

Features

• pytorch 2.3 support
• more gqa group sizes
• add mma instructions for fp8 (#179) (d305798)
• mma rowsum for fp8 (#180) (5af935c)
• support any num_heads for get_alibi_slope (#200) (b217a6f)

Bug Fixes

• fix python package dispatch error message (#182) (8eed01c)