flashinfer-ai/flashinfer v0.7.5rc1

Release v0.7.5rc1

on GitHub

What's Changed

• perf(cake_rmsnorm): blocked FP32 dw sums for the large-T B200 backward windows by @yyihuang in #8848
• fix: Windows build of ln_silu_headers.cuh and Windows import of the b12x package by @ranx2026 in #8300
• docs(msa_ops sm90): list the measured shapes where the Cake route is not faster by @yyihuang in #8852
• perf(cake_kda): round-robin bin arithmetic and a start delay for shorter-bin CTAs in two 128-token turnover builds of BF16 KDA prefill (sm_103a) by @yyihuang in #8874
• fix(quantization): keep the pre-#6207 column loop verbatim for LOADS_IN_FLIGHT=1 by @jimmyzho in #7834
• feat(batch_invariant): mm_mxfp8(batch_invariant=True), the MXFP8 GEMM class by @bkryu in #8536
• perf(cake_dsv4): raw-index SWA V gather in sm_107a FP8 low-head prefill by @yyihuang in #8860
• perf(cake_dsa_indexer_det): split merge reads its staging once into shared 32-bit keys by @yyihuang in #8880
• perf(cake_gdn): sector-complete state I/O in SM107 FP16-I/O and BF16-state gate-stream full-DV prefill (up to 1.21x) by @yyihuang in #8877
• perf(cake_causal_conv1d_det): SM-wide row-group backward for 264 <= C < 296 and 344 <= C < 416 on SM107 by @yyihuang in #8878
• perf(cake_kda_det_kcp): fused forward walk with the chunk output and 8 x 4 affine merge builds, bitwise identical by @yyihuang in #8875
• perf(cake_msa_icp_indexer): TP2 page-256 prefill exchange as a peer publish, two pages per MMA list build by @yyihuang in #8865
• feat(batch_invariant): mm_fp4(batch_invariant=True), the NVFP4 and MXFP4 GEMM classes by @bkryu in #8537
• perf(cake_grouped_fp8_gemm): SM107a lead-prefetch BLOCK_N 256 for padding-heavy short-K layouts by @yyihuang in #8881
• feat(cake_mla_value_gate): add an experimental fused MLA value projection + sigmoid gate by @yyihuang in #8867
• feat(sm100): add opt-in ExpCast for avoiding exponential in S->P for FP8 PV by @adsridhar in #6481
• perf(cake_alpha_moe): NVFP4 fused batch kernel keeps its activation scale in registers on sm_100a by @yyihuang in #8869
• build: write flashinfer/_build_meta.py atomically by @yyihuang in #8883
• perf(cake_dsv4): SM107 FP8 64-head decode takes the M64 body earlier where the four-stage sibling launches by @yyihuang in #8885
• perf(cake_mhc_det): with grad, L2 hints for fn partial sums and launch tuning on sm_100a / sm_107a by @yyihuang in #8870
• perf(cake_dense_projection_det_gemm): drop the 3-split router weight-gradient rule at H 6144 by @yyihuang in #8886
• fix(cake_gdn_det): launch the fwd_ho_vh2 stage on fwd_ho's grid where its record is the fwd_ho program by @yyihuang in #8887
• chore(cake_causal_conv1d_det_train): record the published generator revision by @yyihuang in #8894
• perf(cake_mla): hide the V-shift hand-offs of the NVFP4 MLA paged decode wide route behind loads by @yyihuang in #8884
• fix(cake_kda_det_kcp): CUDA 12.9 floor for the sm_100a 256-bit stores and walk-form provider compatibility by @yyihuang in #8896
• perf(cake_dsa_indexer_det): an FP8 wide scan program on sm_100a by @yyihuang in #8895
• perf(cake_vsa_sm120): overlap softmax with P*V and move mbarrier retries off the block loop by @yyihuang in #8898
• perf(cake_rmsnorm): blocked FP32 dw sums for the large-T GB300 backward windows by @yyihuang in #8899
• feat(cake_mla_value_gate): load the request lengths through L1 by @yyihuang in #8904
• refactor(cake_mhc_det): remove the mHC training names without det by @yyihuang in #8900
• perf(cake_lm_head_loss): two-slice counted dX and a 1946-row zero-gradient floor on sm_107a by @yyihuang in #8902
• fix(cake_kda_train_kcp): CUDA 12.9 floor for the sm_100a PTX ISA 8.8 entries by @yyihuang in #8911
• fix(cake_kda_det): gate-gradient accuracy of the GLM intra stage (diagonal TF32 rounding) by @yyihuang in #8909
• feat(cake_fp8_rowwise_gemm_allreduce): SM90 fused FP8 GEMM + all-reduce for TP2/4/8 by @yyihuang in #8905
• refactor(cake_causal_conv1d_det): det names only, remove the pre-det-rename aliases by @yyihuang in #8912
• perf(cake_alpha_moe): NVFP4 fused batch kernel reduces two output blocks per bulk add on sm_100a by @yyihuang in #8903
• perf(cake_dense_projection_gemm): cluster DSMEM split-K for the thin FP32 weight gradient on 152-SM sm_100a parts by @yyihuang in #8844
• test(cake_gdn): with grad, stop depending on the former no-det names of the deterministic packages by @yyihuang in #8916
• perf(cake_lm_head_loss): cast the weight gradient at the last live chunk of the zero-gradient row skip (sm_100a, sm_103a) by @yyihuang in #8914
• perf(cake_mok_bf16): per-expert x-replay release, expert-major weight-gradient claims and in-place dy pre-weighting in the sm_107a MXFP8 backward (#7821) by @yyihuang in #8919
• perf(moe_ep): faster SM90 MXFP4/NVFP4 pull MegaMoE (N72/N80 token tiles, measured defaults, K128 pipelining) by @ruoqianguo in #7961
• test(cake_gdn_det_kcp): move det KCP tests off the KDA legacy names; refuse other-mode saved sets by @yyihuang in #8918
• test(cake_gdn_train_kcp): force P=2 only on rows whose plan picks P=1 by @yyihuang in #8923
• refactor(cake_pplx): one top-level lazy-export table; Cake backends under experimental/cake keep their sources by @yyihuang in #8932
• perf(cake_kda_det_kcp): leaner host prologue in the KCP forward, bitwise identical by @yyihuang in #8929
• feat(moe_ep): expand the Rubin inference tuning grid by @akaashrp in #8575
• feat(cake_mla_value_gate): add an 8-row tile program for small batches and PDL by @yyihuang in #8917
• fix(cake_kda_det): refuse a non-deterministic workspace or saved set in the det single-rank entries by @yyihuang in #8915
• perf(cake_dsa_indexer_det): an FP8 wide scan program on sm_107a by @yyihuang in #8933
• perf(cake_batch_deepgemm_fp8): 256-byte L2 promotion for three sm_107a programs by @yyihuang in #8938
• perf(cake_vsa_sm120): scale O on every block without a warp vote in vsa_sm120_blk64 by @yyihuang in #8940
• perf(cake_batch_deepgemm_fp8): BLOCK_M 192 route for the sm_103a B=1 M=8192 swap-AB point at expected_m >= 5890 by @yyihuang in #8941
• fix(cake_kda_det_kcp): refuse the non-deterministic package's saved set and workspace with ValueError by @yyihuang in #8943
• refactor(cake_causal_conv1d): rename the non-deterministic package to nondet (#8928) by @yyihuang in #8906
• refactor(cake_gdn_kcp): the GDN KCP facade runs the non-deterministic package only by @yyihuang in #8920
• refactor(cake_moe_grouped_gemm): det-only names, explicit deterministic=False, and deterministic weight-gradient accumulation by @yyihuang in #8924
• chore(cake_kda_kcp): re-render the test-only KCP programs in the current generated form by @yyihuang in #8948
• ci(api-diff): count _LAZY_EXPORTS rows as exports; test: check lazy-export binding before access by @yyihuang in #8937
• docs: add autograd backend naming rules and target table by @yyihuang in #8930
• test(cake): shared check of the mode word in Cake training names (#8928) by @yyihuang in #8947
• refactor(cake_gdn_det_kcp)!: remove the legacy no-det GDN KCP names by @yyihuang in #8921
• refactor(cake_dsa_indexer_det_train): mode-word names for the Cake DSA indexer (#8928) by @yyihuang in #8936
• refactor(cake_kda_det_kcp): remove the pre-det name aliases by @yyihuang in #8901
• refactor(cake_gqa_nondet_train): name the GQA HD256 training package by its mode (#8928) by @yyihuang in #8939
• refactor(cake_rmsnorm_silu_gate): follow #8928 (det/nondet module names, backends under experimental/cake, no deterministic= keyword) by @yyihuang in #8922
• fix(bench_cake_vsa_sm90): give the fa3 arm of a GQA problem its own out layout by @yyihuang in #8958
• perf(cake_dsv4): split SM107 FP8 128-head decode calls of up to 17 query rows across SMs by @yyihuang in #8950
• perf(cake_dsv4): tune SM107 FP8 64-head decode routing for the four-stage sibling by @yyihuang in #8955
• test(cake_package): check that every Cake package imports with flashinfer blocked by @yyihuang in #8951
• perf(cake_rmsnorm_det): compensate every step in the large-token backward on SM100, SM103 and SM107 by @yyihuang in #8835
• ci(pre-commit): cover #8928 autograd backend csrc paths before family renames by @yyihuang in #8957
• perf(cake_mla): trim the V-shift work of the NVFP4 MLA paged decode wide route by @yyihuang in #8956
• test(cake): skip type aliases in the deterministic-keyword check (#8928) by @yyihuang in #8961
• test(cake_gqa_nondet_train): check the shared #8928 name rules with cake_mode_names by @yyihuang in #8962
• refactor(cake_kda_det): remove the single-rank names before the det rename by @yyihuang in #8908
• perf(cake_dsa_indexer_det): pick the FP8 wide unit by wave count on sm_100a by @yyihuang in #8954
• refactor(cake_dsa_sparse_attention_nondet): name the DSA sparse-attention autograd package by its mode (#8928) by @yyihuang in #8942
• perf(cake_deepgemm): route an SM120 batch-128 64K-context decode step to a faster program (RTX 5090 h32 p64 n4) by @yyihuang in #8965
• perf(cake_mega_moe_v3): faster TP4 EP1 and TP8 EP1 T32 / T128 / T512 / T1024 / T4096 rank programs on sm_103a by @yyihuang in #8970
• perf(cake_mxfp8_megamoe_ep16): sm_107a CuTe-DSL fused programs read the peer tables from shared memory by @yyihuang in #8971
• perf(cake_grouped_fp8_gemm): SM107a UE8M0 scale-word reuse for BLOCK_N 128 block-scaled programs by @yyihuang in #8972
• refactor(cake_dense_gemm): rename both packages to det and nondet (#8928) by @yyihuang in #8926
• perf(cake_vsa_sm120): charge a split part more once K and V pass half the L2 by @yyihuang in #8973
• perf(cake_dsa_indexer_det_train): fuse the first radix-select sweep in the sampled calls of the BF16 sm_103a l6 programs by @yyihuang in #8974
• perf(cake_dsv4): split SM107 FP8 128-head decode calls of 14-15 query rows into three pieces by @yyihuang in #8978
• perf(cake_grouped_fp8_gemm): SM107a lead-prefetch route for more BLOCK_N 256 single-run layouts by @yyihuang in #8979
• perf(cake_dsa_indexer_det_train): sm_100a FP8 l6 programs load K_scale one tile ahead by @yyihuang in #8976
• refactor(cake_lm_head_det_train): rename the deterministic LM head training entries and drop the deterministic routing by @yyihuang in #8925
• perf(cake_batch_deepgemm_fp8): two early-setup sm_107a programs for the N=4096 K=7168 and N=7168 K=2048 bands by @yyihuang in #8981
• fix(moe_ep): handle empty local batches in two-sided combine by @lesj0610 in #8821
• perf(cake_gdn): prefetch the initial state into L2 in sm_107a gate-stream prefill (up to 1.12x) by @yyihuang in #8982
• feat(experimental): add SM100/SM103 HC sum-state backend by @StudyingShao in #8944
• perf(cake_dsa_indexer_det_train): sm_107a FP8 runs an l6 snake program on short single-document calls by @yyihuang in #8986
• feat(trace_apply): dynamic loading of verified kernel packages by @akshaver in #8866
• feat(prims-ts): support independently inferred device global scales by @yuxianq in #8093
• perf(cake_mok_bf16): stream macro batch 0's unaliased minibatches ahead of macro batch 1's combine in the sm_107a MXFP8 forward (#7821) by @yyihuang in #8988
• perf(cake_dsa_indexer_det_train): take the wide unit from 5.5 rounds per CTA on sm_100a by @yyihuang in #8989
• ci(docker): install rsync in the CI images by @dierksen in #8990
• perf(cake_kda): enable the three 128-token turnover builds of BF16 KDA prefill on sm_100a by @yyihuang in #8994
• perf(cake_mla): one relaxed exit rendezvous for the NVFP4 MLA paged decode wide route by @yyihuang in #8977
• feat(moe_ep): add isolated Kimi K3-family MegaMoE backend by @Anerudhan in #8934
• perf(cake_dsv4): early-PV sm_107a FP8 low-head prefill for launches up to 64 query tokens by @yyihuang in #8991
• feat(cake_mega_moe_v3): fence the in-launch global hand-offs before async-proxy reads by @yyihuang in #9001
• perf(cake_vsa_sm120): merge a two-part row with one partial read in vsa_sm120_blk64 by @yyihuang in #9002
• fix(mamba): bound FP8 ReplaySSM checkpoint scales by @askliar in #8968
• chore(cake): prepare LM-head nondet pre-commit and fix GDN device scoping (#8928) by @yyihuang in #8985
• perf(cake_dsv4): split SM107 FP8 64-head decode calls of few query rows by @yyihuang in #8984
• fix: support Windows file locking and skip unsupported Quack tests by @saltyminty in #8714
• refactor(cake_moe_grouped_gemm): name the MoE grouped GEMM modules by their mode (#8928) by @yyihuang in #8963
• refactor(cake_causal_conv1d): rename deterministic autograd APIs and relocate backend by @yyihuang in #8953
• feat: support token-strided QKV in CuTe KDA prefill by @syuoni in #8613
• test(cake_dense_gemm): run the shared #8928 mode-name checks in both dense names tests by @yyihuang in #9010
• docs: release note for the #8928 CAKE autograd renames by @yyihuang in #8960
• Add opt-in Rubin CuTe DSL prefill backend by @saltyminty in #8716
• perf(cake_kda): preload the next task's entering state into tensor memory in a new work-bin build of BF16 KDA prefill (sm_103a) by @yyihuang in #9012
• test(cake_kcp): always-on CPU parity of the KCP partition and base context across the four packages (#8928) by @yyihuang in #9011
• perf(cake_fmha): BF16 paged-context runtime-size programs for the non-causal, LSE and sink forms by @yyihuang in #9016
• refactor(cake_kda_det)!: name the det single-rank KDA entries by mode and move its backend under experimental/cake (#8928) by @yyihuang in #9008
• perf(cake_dsa_indexer_det_train): walk key tiles in descending order in the sampling-off sm_103a wide programs by @yyihuang in #9019
• perf(cake_rmsnorm_silu_gate): route the deterministic forward by rows per SM by @yyihuang in #9007
• refactor(cake_mhc): name the mHC training modules by their mode: cake_mhc_det_train, cake_mhc_nondet_train, cake_mhc_head_det_train (#8928) by @yyihuang in #8907
• refactor(cake_rmsnorm): mode-word names and experimental/cake backends (#8928) by @yyihuang in #8910
• refactor(cake_lm_head_nondet_train): rename the non-deterministic LM head training package (#8928) by @yyihuang in #8952
• docs(cake_moe_grouped_gemm): state which launches allocate stream workspaces by @yyihuang in #9022
• perf(cake_deepgemm): route an SM120 batch-128 64K-context decode step to a faster program (RTX 5090 h64 p64 n3) by @yyihuang in #9023
• test(comm): tolerate reference rounding in MXFP8 norm-output cases by @righthandabacus in #8997
• docs(training_backend_migration): add the 0.7.4rc1 det names of cake_rmsnorm_train_backend by @yyihuang in #9045
• fix(cake_rmsnorm_silu_gate): refuse other-mode workspace on empty input by @yyihuang in #9026
• fix(gdn): order TMEM dealloc after the compute warps' tcgen05.ld in the Blackwell GDN kernels by @bastianhagedorn in #8854
• perf(cake_rmsnorm_silu_gate_nondet_train): FP32-fold sm_107a backward for 1024 to 3200 rows per SM by @yyihuang in #9015
• perf(cake_dense_gemm_det_train): restore the 3-split router weight-gradient rule at H 6144 by @yyihuang in #9034
• perf(cake_causal_conv1d): 32-vector channel tiles for the non-deterministic backward at C224 and C520-760 on sm_100a / sm_103a by @yyihuang in #9035
• refactor(cake_gdn): rename the GDN autograd modules to det / nondet (#8928) by @yyihuang in #8959
• perf(cake_dsa_indexer_det_train): sm_103a FP8 runs one-document wide calls in the :lsf copy by @yyihuang in #9021
• feat(batch_invariant): batch-invariant attention behind batch_invariant=True on the paged decode, paged prefill and ragged prefill wrappers by @Vinnie6167 in #8590
• feat(cake_kda): with grad, in-kernel Q/K L2 norm in the non-deterministic single-rank and context-parallel packages by @yyihuang in #8801
• perf(cake_kda): owner/helper hybrid build of the BF16 KDA prefill M64 module for equal-length 2048+ token launches on GB300 (sm_103a) by @yyihuang in #9040
• perf(cake_mhc_det): with grad, L2 hints for fn partial sums above 8192 tokens and mapping-pass tuning on sm_100a by @yyihuang in #9042
• perf(cake_dense_gemm_det_train): sm_107a three-chain fp32 weight gradients and ragged weight-gradient rules by @yyihuang in #9043
• fix(cake_kda): transpose the BT16 chunk inverse as a 16x16 matrix and compose its off-diagonal block from the complete 8x8 inverses by @burak-inherent in #6132
• fix(cake_dsv4): run the SM107 FP8 64-head ragged 2x257 prefill row on the M64 body by @yyihuang in #9027
• perf(cake_vsa_sm120): split packed plans less once K and V pass 48 MB by @yyihuang in #9046
• refactor(cake_kda_nondet)!: name the non-deterministic KDA training modules by mode, with the det KCP rename (#8928) by @yyihuang in #9017
• perf(cake_dsa_indexer_det_train): sm_103a FP8 picks the wide unit by waves of the grid by @yyihuang in #9047
• fix(kda): isolate non-finite V padding in chain prefill by @michaelfeil in #8599
• perf(cake_mok_bf16): two-phase all-thread quantize and next-task prefetch in the sm_107a MXFP8 backward dispatch (#7821) by @yyihuang in #9050
• feat(cake_mhc_head_nondet_train): add the non-deterministic mHC training head module (#8928) by @yyihuang in #9044
• perf(trtllm-gen): opt-in wave-fill KV split for decode (FLASHINFER_TRTLLM_GEN_MULTI_WAVE_KV_SPLIT=2) by @shyeh25 in #8744
• perf(cake_gdn): sm_107a gate-stream GDN prefill for FP16 I/O and BF16-state Hq2/Hv8 (up to 1.28x) by @yyihuang in #9053
• perf(cake_fmha): interleaved tail order for the BF16 runtime-size sink and LSE programs by @yyihuang in #9055
• perf(cake_dense_gemm_nondet_train): cluster split-K FP32 weight gradient on 152-SM GB300 (sm_103a) and 148-SM B200 by @yyihuang in #9051
• perf(cake_gdn_nondet_train): split the fused forward scan and output over the value halves on R200 by @yyihuang in #9056
• perf(cake_kda_det): roll the sm_107a token-row drains of the det GLM intra kernels (follow-up to #7817) by @yyihuang in #9048
• perf(cake_kimi_k3_fp8_projection): SM100a grouped tile raster with two TMEM accumulators for TP8 q_proj M 4097..16384 by @yyihuang in #9059
• refactor(cake_causal_conv1d_nondet_train): name the host-plan cache helper by its mode (#8928) by @yyihuang in #9057
• refactor(cake_gqa_nondet_train): keep default_workspace internal to cake_backend (#8928) by @yyihuang in #9054
• perf(cake_dsa_indexer_det_train): plain l6 order for packed and few-wave sampling-off calls on sm_107a by @yyihuang in #9052
• perf(cake_mega_moe): sm_107a fence-free push combine for the BF16 rank-major EP8 layer by @yyihuang in #9061
• perf(cake_dense_gemm_nondet_train): exact-FP32 cluster SIMT router forward at 1 to 192 tokens on 212-SM sm_107a by @yyihuang in #9060
• perf(cake_fmha): sm_103a FP8 paged-context form programs with the 4-stage K/V ring by @yyihuang in #9062
• perf(cake_kda): enable the tensor-memory next-state preload work-bin build of BF16 KDA prefill on B200 (sm_100a) by @yyihuang in #9063
• fix(cake_kda_kcp): workspace bytes use the context device's architecture; strict #8928 names in the KCP parity test by @yyihuang in #9058
• perf(cake_kimi_k3_mla): persistent two-CTA wide decode route on B200 by @yyihuang in #9066
• perf(cake_alpha_moe): NVFP4 fused batch kernel reduces two output blocks per bulk add on sm_103a by @yyihuang in #9064
• perf(cake_dense_gemm_nondet_train): router weight-gradient head/tail windows for 152-SM GB200 and GB300 by @yyihuang in #9069
• perf(cake_lm_head_det_train): skip zero-weight policy rows and zero-dlogp log-probability rows by @yyihuang in #9071
• perf(cake_dense_gemm_det_train): swap single-head weight gradients with a narrow side in the generic entry by @yyihuang in #9070
• perf(cake_dsa_indexer_det_train): sm_103a FP8 lets short calls take the wide unit from fewer tiles by @yyihuang in #9065
• fix(cake_mega_moe_v3): proxy fence in the carried programs, coherent loads of in-launch data by @yyihuang in #9074
• feat(cake_kimi_k3_mla): add LSE, DCP and non-causal variants to the Cake Kimi-K3 MLA decode kernel by @yyihuang in #9067
• perf(cake_kda): preload the next task's entering state into tensor memory in the N16 work-bin build of BF16 KDA prefill (sm_103a) by @yyihuang in #9075
• feat(batch_invariant): softmax(batch_invariant=True), sampling.log_softmax and sampling.log_probs: the row log-sum-exp class lse.row by @bkryu in #8999
• perf(cake_dsa_sparse_attention_nondet): backward register budgets on SM100a, SM103a and SM107a by @yyihuang in #9076
• perf(cake_fmha): sm_103a FP8 paged-context one-stage form programs by @yyihuang in #9073
• docs(training_backend_names): allow a KCP route inside its single-rank backend package by @yyihuang in #9082
• perf(cake_kda_det_kcp): run the forward's summary pair on two streams when K1 fills a wave by @yyihuang in #9072
• perf(cake_dsa_indexer_det_train): split calls fuse the slab rank finalize into the merge by @yyihuang in #9078
• perf(cake_dsa_indexer_det_train): split count from the merge-resident limit for few-unit calls on sm_100a and sm_103a by @yyihuang in #9081
• perf(cake_alpha_moe): NVFP4 aligned entry drains a zero-at-rest FP32 accumulator instead of seeding it by @yyihuang in #9080
• perf(cake_gdn_nondet_train_kcp): v2 forward output walk with a re-ordered vh 64 chunk step by @yyihuang in #9083
• perf(cake_batch_deepgemm_fp8): route single-wave sm_107a N=4096 K=7168 keys with B >= 2 to the early-setup BLOCK_N 96 program by @yyihuang in #9084
• perf(cake_rmsnorm_nondet_train): sm_103a and sm_107a small-T backward runs the det programs by @yyihuang in #9079
• perf(cake_gdn): route ragged sm_107a BF16-state, FP16-I/O and no-initial-state prefill to DV-split (up to 1.18x) by @yyihuang in #9085
• perf(cake_msa): blackwell_msa paged FP8-Q prefill (146 query tokens, 64/4 heads, 2 pages) on sm_103a with the existing union program by @yyihuang in #9086
• fix(cake_dsv4): reject sparse index tables the index-row loads cannot read by @yyihuang in #9088
• perf(cake_fmha): route more BF16 sink/LSE and non-causal pack-16 calls to the AOT context programs by @yyihuang in #9090
• perf(cake_mega_moe): sm_100a PB, G3, CAB4 and push combine for the BF16 rank-major EP8 layer by @yyihuang in #9091
• perf(cake_batch_deepgemm_fp8): sm_107a scale prefetch for the BLOCK_N 32 band and early setup for one-wave BLOCK_N 192 keys by @yyihuang in #9092
• perf(cake_dsv4): three-piece split merge for sm_107a FP8 decode at 16-17 rows by @yyihuang in #9089
• test(gemm): validate architecture-specific grouped FP8 routes by @righthandabacus in #9039
• test(norm): accept memset zero fills in captured backward by @righthandabacus in #8996
• fix(cake_lm_head_det_train): keep the zero-weight split within the plain call's peak memory by @yyihuang in #9098
• perf(cake_gdn): add sm_107a full-DV GDN prefill twins with sector-complete state I/O (up to 1.17x) by @yyihuang in #9099
• refactor(cake_gated_act_mxfp8): follow #8928 (det module name, backend under experimental/cake) by @yyihuang in #9100
• perf(cake_mega_moe_v3): shorter expert lookup and a parallel completion signal in the E384 T4 sm_103a programs by @yyihuang in #9102
• perf(cake_fp8_rowwise_gemm_allreduce): per-call reducer layout for SM90 TP2/4/8 by @yyihuang in #9101
• perf(cake_gdn_det_train): read q and k in place from packed qkv by @yyihuang in #9103
• fix(cake_lm_head_det_train): keep the backend and a signed zero dW on calls without a gradient GEMM by @yyihuang in #9105
• fix(cake_kda_nondet): make the checkpoint hand-off switch keyword only (#7546) by @yyihuang in #9109
• fix(cake_alpha_moe): harden the NVFP4 aligned accumulator cache (follow-up to #9080) by @yyihuang in #9110
• perf(cake_gdn_det_train): read v in place from packed qkv by @yyihuang in #9115
• feat(cake_dsa_indexer_nondet_train): set-repeatable unordered DSA indexer top-k for training by @yyihuang in #9113
• perf(cake_causal_conv1d_det_train): more 32-vector tiles in backward and BF16-tap forward by @yyihuang in #9106
• cake_gdn_det_train: price the backward adjoint table for the tensor-core scan by @yyihuang in #9118
• perf(cake_gated_act_mxfp8): 16-lane rowwise program for small calls on SM100 and SM107 by @yyihuang in #9120
• perf(cake_gqa_nondet_train): faster split backward; every output repeats bit for bit by @yyihuang in #9116
• perf(cake_mla): leaner softmax, transform and correction work in the NVFP4 MLA paged decode wide route by @yyihuang in #9114
• fix(cake_causal_conv1d_det_kcp): non-cooperative KCP backward by default, no environment read by @yyihuang in #9111
• perf(cake_batch_deepgemm_fp8): sm_107a builds of the packed-UE8M0 swap-AB M32 serving GEMM by @yyihuang in #9127
• docs(cake_dsa_indexer_det_train): per-output repeat statement and doc fixes by @yyihuang in #9125
• perf(cake_dsv4): TMEM-batch SM107 FP8 64-head decode program for MTP launches of 108+ rows by @yyihuang in #9128
• perf(cake_causal_conv1d_det_train): sm_107a backward runs the stage-pair program for C 520-1016 by @yyihuang in #9126
• perf(cake_fmha): route more sm_100a BF16 causal SINK/LSE calls to the AOT context programs by @yyihuang in #9131
• perf(cake_lm_head_det_train): skip the logits GEMMs in the checkpoint recomputation of the log-probability entry by @yyihuang in #9124
• perf(cake_dsv4): per-count split merges for sm_107a FP8 decode at 1-8 rows by @yyihuang in #9136
• perf(cake_dsa_indexer_nondet_train): sm_107a ticket scan programs and sampling rules by @yyihuang in #9130
• feat(cake_gqa_det_train): deterministic packed-varlen HD256 GQA training attention by @yyihuang in #9121
• fix(cake_mhc): fold dfn partials in chunk order so the non-deterministic backward repeats bit for bit by @yyihuang in #9123
• perf(cake_gqa_nondet_train): side-stream bwd plan, exact chunks from host seqlens by @yyihuang in #9140
• perf(cake_mok_bf16): default the communication-SM split to the measured rows by @yyihuang in #9107
• perf(cake_rmsnorm_silu_gate_det_train): one-kernel backward and persistent forward by @yyihuang in #9143
• perf(cake_dsa_indexer_det_train): sm_107a sampled-threshold rules for packed and long-document calls by @yyihuang in #9135
• perf(cake_deepgemm): route four SM120 RTX PRO 6000 h64 p64 decode steps to faster programs by @yyihuang in #9150
• fix(cake_dsa_indexer_nondet_train): CTA-pair ticket programs wait for the last echo by @yyihuang in #9147
• perf(cake_causal_conv1d_det_train): sm_107a forward token floor and C 2048 channel-band programs by @yyihuang in #9133
• perf(cake_gqa_det_train): side-stream bwd plan, exact chunks from host seqlens by @yyihuang in #9151
• fix(cake_gqa_nondet_train): host-boundary checks at every size, exact range check past 2^31 tiles, nine-argument apply by @yyihuang in #9149
• fix(cake_dsv4): bound source-exact decode rows by num_query_tokens by @yyihuang in #9156
• perf(cake_dsa_indexer_det_train): sm_107a l6 and CTA-pair l6 draw their work units from tickets by @yyihuang in #9153
• perf(cake_fmha): sm_100a one-stage FP8 context form programs by @yyihuang in #9160
• perf(cake_rmsnorm_det_train): faster backward routes for residual H 6144, H 2048, H 4096 and FP32-weight H 5120 on SM100, SM103 and SM107 by @yyihuang in #9134
• perf(cake_moe_grouped_gemm_det_train): fused FP32 wgrad accumulate, up to 1.25x faster by @yyihuang in #9132
• perf(cake_mhc_det): SM103 fwd keeps L2-hinted X read to 6656 (was 6144) by @yyihuang in #9157
• perf(cake_causal_conv1d_det_kcp): packed FP32 backward, one-SM-short forward by @yyihuang in #9142
• perf(cake_dsa_sparse_attention_nondet_train): backward wait hint on SM100 / SM103 and a per-architecture NoPE dKV window by @yyihuang in #9163
• perf(cake_mhc_head_det_train): per-H 64-bit SM100/103 kernels, SM107 L2 bulk-reduce acc bwd by @yyihuang in #9162
• perf(cake_kda_det_train_kcp): faster backward and checkpointed step on the deterministic KDA KCP route by @yyihuang in #9141
• perf(cake_dsv4): split sm_107a FP8 128-head decode at 18-20 rows by @yyihuang in #9167
• perf(cake_gated_act_mxfp8): hybrid band path and per-size band-path routing on SM100, SM103 and SM107 by @yyihuang in #9138
• cake_gdn_det_train: copy unaligned q/k/v/do to an aligned buffer before launch by @yyihuang in #9169
• fix(cake_dsa_indexer_det_train): CTA-pair ticket programs wait for the last echo by @yyihuang in #9166
• perf(cake_dense_gemm_nondet_train): narrow-output programs and accurate split-K weight gradients by @yyihuang in #9171
• perf(cake_gqa_nondet_train): run the host-boundary checks on side streams beside bwd_pre by @yyihuang in #9165
• fix(cake_moe_grouped_gemm_det_train): address review comments from #9132 by @yyihuang in #9173
• fix(cake_lm_head_det_train): keep the default step's peak within the logits-materializing bound by @yyihuang in #9152
• perf(cake_fp8_rowwise_gemm_allreduce): grid lockstep throttle for large SM90 calls by @yyihuang in #9164
• perf(cake_kda_nondet_train_kcp): faster backward and checkpointed step on the non-deterministic KDA KCP route by @yyihuang in #9155
• perf(cake_fmha): sm_103a NVFP4 context programs release the scale slot early and deepen the packed K/V ring by @yyihuang in #9177
• fix(cake_lm_head_nondet_train): no FP32 dW accumulator, the call's backend and dW = cast(0 * g) on calls without a gradient GEMM by @yyihuang in #9179
• perf(cake_batch_deepgemm_fp8): faster masked_m setup on sm_107a for N=7168, K=2048 by @yyihuang in #9183
• perf(cake_kda_det_train_kcp): replay the backward tapes on a side stream by @yyihuang in #9168
• fix(cake_gdn): end the tile-boundary deadlock in sm_107a initial-state prefill by @yyihuang in #9184
• test(cake_dsa_indexer_nondet_train): match the deterministic sm_107a ticket programs after #9153 by @yyihuang in #9185
• fix(cake_gdn): end the tile-boundary deadlock in sm_100a and sm_103a initial-state prefill by @yyihuang in #9187
• perf(cake_rmsnorm_silu_gate_det_train): approximate-gate backward by @yyihuang in #9188
• perf(cake_gqa_det_train): run the host-boundary checks on side streams beside bwd_pre by @yyihuang in #9180
• perf(cake_mega_moe_v3): deeper operand stages for the TP8 EP8 4-token rank program by @yyihuang in #9189
• cake_gdn_det_train: read cached device tables after their write on any stream; captures take their own workspaces by @yyihuang in #9186
• feat(cake_dsa_sparse_attention_det_train): deterministic DSA sparse MLA training attention by @yyihuang in #9161
• perf(batch_invariant): skip backend selection for the batch_invariant=True GEMM APIs; cache device checks by @qiching in #8598
• perf(cake_moe_finalize_allreduce_fusion): run the large-token shared program from 896 tokens by @yyihuang in #9192
• feat(cake_mhc_det): serve hidden 5120 / 7168 on SM107 by @yyihuang in #9181
• feat(cake_moe_grouped_gemm_nondet_train): run-to-run reproducible k512 split tail and evict-first routes by @yyihuang in #9139
• perf(cake_dsv4): SM107 FP8 64-head few-row decode: TMEM-batch one-token window and last-tile L2 prefetch program by @yyihuang in #9196
• perf(cake_gated_act_mxfp8): retuned per-size programs on SM103 and SM107 by @yyihuang in #9195
• perf(cake_gdn_nondet_train): split forward on long documents, A-form forward chain and one-wave backward split by @yyihuang in #9175
• perf(cake_causal_conv1d_det_kcp): overlapped KCP backward, rolled sm_107a backward by @yyihuang in #9190
• perf(cake_batch_deepgemm_fp8): faster masked_m setup for N=512, K=128 on sm_100a/sm_103a by @yyihuang in #9199
• docs(cake_mhc_nondet): state the SM107 hidden-size exception in the registry text by @yyihuang in #9200
• perf(cake_causal_conv1d_nondet_kcp): overlapped KCP backward that claims its work at run time by @yyihuang in #9154
• perf(cake_gdn): cache nontranspose decode launch, TMA-state body for sm_107a B10-20, up to 1.37x by @yyihuang in #9202
• fix(cake_fmha): paged context programs ignore NaN in V rows past the sequence end by @yyihuang in #9203
• perf(cake_dense_gemm_det_train): MLA batch raster, serpentine K walk and exact-T rules for the projection GEMM by @yyihuang in #9170
• perf(cake_gqa_det_train): forward tiles of two query heads, 32-byte O stores by @yyihuang in #9205
• perf(cake_rmsnorm_nondet_train): faster backward with pair-row kernels and accuracy-fixed routes on SM100, SM103 and SM107 by @yyihuang in #9178
• perf(cake_dense_gemm_nondet_train): in-cluster split-K for thin weight gradients and repeatable default plans by @yyihuang in #9194
• perf(cake_dsa_indexer_det_train): sm_107a persistent rank finalize spares one SM for another stream by @yyihuang in #9207
• fix(cake_gdn_nondet_train): cached device tables across streams and CUDA graph captures by @yyihuang in #9210
• perf(cake_dsv4): run sm_107a FP8 low-head prefill launches of up to 27 rows on a V-block sibling by @yyihuang in #9201
• fix(cake_kda_det_train_kcp): no extra reserved memory from kept checkpoint rows by @yyihuang in #9206
• perf(cake_causal_conv1d_det_train): even channel tiles for the C 576 / 640 forward from per-arch token counts by @yyihuang in #9214
• perf(cake_mega_moe_v3): run the TP4 EP1 1-token rank route on a 148-CTA persistent grid by @yyihuang in #9216
• feat(prims_ts): fp8 [32,32] block-scale routed MoE on Prims-TS by @feih-nv in #8666
• perf(cake_gqa_nondet_train): forward tiles of two query heads, 32-byte O stores by @yyihuang in #9208
• perf(cake_dsa_sparse_attention_nondet_train): lighter host path for the RoPE varlen forward and the NoPE entries by @yyihuang in #9213
• perf(attention): group tiny-KV FP8 context for uniform Q1-3 by @saltyminty in #9018
• perf(cake_gdn_det_train): forward on the WY stage's A and split forward on long documents by @yyihuang in #9204
• fix(cake_kda_nondet_train_kcp): keep the checkpoint hand-off's kept set in a private memory pool by @yyihuang in #9193
• perf(cake_fmha): sm_103a NVFP4 paged context drops a redundant proxy fence by @yyihuang in #9228
• perf(cake_dsa_indexer_det_train, cake_dsa_indexer_nondet_train): sm_107a l6 snake ticket program reads its next ticket at the end of each unit by @yyihuang in #9218
• perf(cake_gated_act_mxfp8): retune the per-size programs on SM100, SM103 and SM107 by @yyihuang in #9226
• perf(cake_causal_conv1d): faster nondet KCP backward host flow by @yyihuang in #9230
• perf(cake_rmsnorm_silu_gate_det_train): look-ahead backward for long packed rows by @yyihuang in #9223
• perf(cake_causal_conv1d_det_train): SM-wide row-group backward for C 520-615 from per-arch token floors by @yyihuang in #9240
• fix(cake_dense_gemm_nondet_train): keep the serpentine K walk out of program knobs by @yyihuang in #9241
• perf(cake_fmha): faster FP8 context LSE, sink and non-causal forms on B200 and GB300 by @yyihuang in #9231
• perf(cake_fmha): add BF16 paged-context sink and LSE programs without tail-order code by @yyihuang in #9236
• perf(cake_kimi_k3_latent_moe): faster Kimi-K3 LatentMoE TP8 decode tail on GB200 and GB300 by @yyihuang in #9176
• perf(cake_mhc_nondet): SM107 bwd fn split waits for the token-owner pass in its last CTA only by @yyihuang in #9235
• fix(cake_hopper_msa): sm_90a top-k keeps saturated-floor and NaN scores selectable by @yyihuang in #9243
• perf(cake_rmsnorm_silu_gate_nondet): faster non-deterministic impl by @yyihuang in #9182
• perf(cake_rmsnorm_det_train): pipelined backward routes for H 2048, H 4096 and H 6144 on SM100 and SM103 by @yyihuang in #9217
• perf(cake_dsa_sparse_attention_nondet): launch before host work on the NoPE backward and RoPE varlen forward; key the NoPE backward by its saved rows by @yyihuang in #9227
• perf(cake_gdn_det): faster backward for value-head groups of three (HV = 3H) by @yyihuang in #9238
• perf(cake_causal_conv1d_det_kcp): single-rank forward program as the KCP forward's main launch by @yyihuang in #9225
• perf(cake_lm_head_det): split-accumulator budget form and device row count on multi-chunk calls by @yyihuang in #9211
• perf(cake_mok_bf16): faster BF16 forward and recompute, context-free forward, register-path epilogues by @yyihuang in #9191
• fix(cake_dsa_sparse_attention_det): key the NoPE backward by its saved tensors, check the varlen cut by @yyihuang in #9239
• perf(cake_fmha): sm_103a NVFP4 paged context dequantizes V on idle warps by @yyihuang in #9253
• perf(cake_fp8_rowwise_gemm_allreduce): drain all-gather from every peer on SM90 by @yyihuang in #9254
• perf(cake_gdn_nondet_kcp): exact-zero adjoint walk cut and light_akk default save mode by @yyihuang in #9174
• perf(cake_gdn_det_kcp): faster, and checkpointed step of the deterministic GDN KCP package by @yyihuang in #9172
• perf(cake_mhc_head_det): SM100 hidden-4096 in 384-thread CTAs by @yyihuang in #9257
• perf(cake_lm_head_nondet): ordered dX slab, one-pass one-row CE, memory budget, kept and vocabulary-major backward by @yyihuang in #9259
• perf(cake_batch_deepgemm_fp8): masked_m-once BLOCK_N 32 build for N=512, K=128 with at most 8 groups on sm_100a/sm_103a by @yyihuang in #9263
• perf(cake_mega_moe_v3): K-256 FC1 stages for TP8 EP1 at 4 tokens; 148-CTA grid for TP4 EP4 and TP8 EP8 at 1 token by @yyihuang in #9264
• perf(cake_dsv4): spread the masked sliding-window rows of sm_107a FP8 low-head prefill over the L2 by @yyihuang in #9262
• fix(cake_hopper_msa): sm_90a top-k keeps zero and denormal scores by @yyihuang in #9265
• perf(cake_causal_conv1d): faster nondet + KCP by @yyihuang in #9266
• perf(cake_dsa_sparse_attention): faster det and nondet bwd on SM100a, SM103a and SM107a by @yyihuang in #9267
• perf(cake_dsa_indexer_det & nondet): sm_107a l6 ticket programs hand out packed key segments longest-first by @yyihuang in #9268
• perf(cake_gated_act_mxfp8): retune the per-size programs on SM100 and SM107 by @yyihuang in #9269
• perf(cake_rmsnorm_silu_gate_nondet): look-ahead backward on long packed inputs by @yyihuang in #9270
• perf(cake_gqa_nondet_train, cake_gqa_det_train): the dQ GEMM writes BF16 dQ; no dQ accumulator clear, no dQ postprocess by @yyihuang in #9271
• perf(cake_rmsnorm_nondet_train): faster backward windows on SM100, SM103 and SM107 by @yyihuang in #9273
• fix(gemm): key cuBLASLt tactics on algorithm descriptors, not heuristic indices by @Vinnie6167 in #6148
• fix(moe): accept strided partitioned C views in the Prims-TS BF16 launch IO by @jimmyzho in #9036
• fix(cake_kimi_k3_latent_moe): cluster-scope cross-CTA landing barriers; GB300 ring-rows tail and full-tile front routes by @yyihuang in #9275
• fix(benchmarks): enable unified NVFP4 MoE on SM107 by @lunarz-dev in #9220
• fix(ci): retry sccache release downloads by @cindyzxq in #8931
• test(batch_invariant): retry the profiled launch when the trace drops its kernel record by @jimmyzho in #9030
• feat(moe_ep): autotune complete Mega and split EP plans by @akaashrp in #8576
• perf(cake_lm_head_det_train): faster budget forms, log-prob calls and sm_107a GEMMs by @yyihuang in #9272
• perf(cake_rmsnorm_det_train): column-task backward programs at small token counts and new small-T forward programs by @yyihuang in #9277
• feat(mla): Declare and test compact-query sparse TRTLLM-GEN dimensions by @dbari in #6249
• ci: bound AI scope evidence before requesting recommendations by @righthandabacus in #9024
• perf(cake_moe_finalize_allreduce_fusion): two-shot shared programs at world sizes 4 and 8 by @yyihuang in #9284
• perf(trtllm-gen): require longer mainloops for multi-wave KV splits by @PerkzZheng in #9261
• perf(prims_ts): restore runtime QToken sparse attention performance by @PerkzZheng in #9222
• fix(kda): fence Small-BH K1 gate-stage writes for the async proxy before their TMA store by @longlee0622 in #9229
• perf(cake_moe_grouped_gemm_nondet_train): queue-walk weight-gradient tiles on SM100 and SM103 by @yyihuang in #9287
• feat(gdn): accept_paths for tree commits; tracked commits on the tcgen05 fold by @kahyunnam in #6257
• feat(gdn): parent-indexed tree verify on the ReplaySSM tcgen05 kernel, with an fp64 reference by @kahyunnam in #6155
• test: reorganize tests/cpu and tests/test_helpers by op so op owners can approve their tests by @bkryu in #9006
• feat(kda): fp32 state for the CuTe recurrent decode kernels by @kahyunnam in #8833
• docs: fix cake references, grouped GEMM formulas, and alias checks by @kangbintNV in #9215
• feat(prims_ts): enable DeepSeek fp8 32x32 scale blocks on SM107 by @feih-nv in #9221
• fix(moe): read the activation-matrix doc as UTF-8 by @feih-nv in #9237
• test: fix CPU tests broken by #9155, #8666 and #8590 by @dierksen in #9291
• Update Docker CI tags to 20261106-54b9f685 by @flashinfer-bot in #6266
• fix(quantization): route NVFP4 quantize to the TMA kernel from m=2048 on SM107 by @kahyunnam in #9292
• perf(cake_mhc_det): backward forms that share an SM with a side kernel on SM100 / SM103 / SM107 by @yyihuang in #9297
• perf(cake_dense_gemm_nondet_train): faster BF16 dense GEMM routes on B200, GB200, GB300 and R200 by @yyihuang in #9299
• feat(gemm): Lamport peer-store epilogue for the low-M BF16 GEMMs by @wookjeHan in #8945
• fix(quantization): reject unsupported CUDA MXFP8 devices by @YaoleWu in #7986
• test(cake_ssd_combined): mock program loading in the Rubin twins sour… by @nvamyt in #8297
• fix(sampling): serialize shared generator RNG reservations by @tlysanhuo in #9148
• feat(batch_invariant): gemm_fp8_nt_groupwise(batch_invariant=True), the block-scale fp8 GEMM class by @bkryu in #9280
• perf(cake_comm): SM90 decoupled-storer producer and SM90 / SM100 / SM103 plan updates for QKV norm-RoPE exchange by @yyihuang in #9276
• fix(moe_ep): correct standard Blackwell NVFP4 configuration and tuning by @akaashrp in #9245
• test(mhc): remove gitignore dependency from package wiring checks by @righthandabacus in #9302
• fix(moe_ep): validate Kimi staged outputs and async flag batches by @akaashrp in #9197
• test(cake_rmsnorm): count a det-pair window as det-inherited when it runs any det program of its route by @jimmyzho in #9038
• fix(moe): exclude unsupported NVFP4 LoRA tile 192 by @jiahanc in #9212
• perf(moe): reduce the host overhead of the SM90 CuTe DSL fused MoE by @Aneureka in #9246
• perf(cake_deepgemm): route eight SM120 RTX 5090 paged MQA decode steps to faster programs by @yyihuang in #9313
• fix(prims-ts): work around QuACK AST rewrite in MLA reduction by @yuxianq in #9087
• fix(cake_rmsnorm): cap the H 2048 backward main kernels at 120 registers so they start beside an all-SM kernel by @jimmyzho in #9037
• test(comm): use a private file store for MXFP8 rendezvous by @righthandabacus in #8998
• feat(moe): add grouped packed NVFP4 Tensor Core candidates to Frost by @YangXu1990uiuc in #7384
• bump version to 0.7.4 by @jimmyzho in #8876
• perf(gemm): cache and parallelize MXFP8 CuTe-DSL autotune compilation by @YangXu1990uiuc in #6615
• feat(deepseek_v41): experimental Prims-TS decode backend for DS4.1 sparse MLA (SM100/SM103/SM107) by @kris1025 in #6988
• feat(gemm): add mm_bf16_fp8 / mm_fp16_fp8, a W8A16 FP8 GEMM for Hopper and Blackwell (CuTe-DSL) by @ishovkun in #7495
• perf(moe_ep): select MXFP4 defaults by expert fill and FC1 depth by @achartier in #9033
• fix(batch_invariant): sampling.log_probs / sampling.log_softmax default to batch_invariant=False. by @bkryu in #9295
• fix(build): retain HC JIT sources in release wheels by @StudyingShao in #9244
• perf(gemm): mm_bf16 pdl=True autotunes every backend; pdl=None now the same as pdl=False by @arpera in #8414
• fix(attention): acquire at cluster scope on the Rubin MTP correction-factor wait by @bastianhagedorn in #8856
• perf(sampling): accumulate deterministic AIR top-p bins with native 32-bit atomics by @ZenAlexa in #8913
• KDA for KDA updated: 3.512x speedup by @smoothsmooth in #9247
• perf(norm): keep 32-bit row offsets in CUDA LayerNorm when they fit by @yuhengcai1 in #8935
• fix(topk): order three cross-CTA group-state accesses in the multi-CTA radix kernels by @yekerr in #8832
• fix(rope): guard rotary_dim <= head_dim to prevent silent cross-head corruption by @SuperMarioYL in #4349
• fix(rope): make cos/sin cache RoPE correct for any rotary_dim and layout by @0z5a in #9104
• test(jit): skip Linux CI bash cases on Windows; read repo files as UTF-8 by @dierksen in #9288
• ci(aot): cache PR AOT nvcc compiles with the rapidsai/sccache fork (host C++ uncached) by @dierksen in #9031
• ci: build the jit-cache shim wheel on the runner, not in the builder image by @dierksen in #8983
• ci(release-build): build flashinfer-cubin with CPU torch to stop runner disk exhaustion by @dierksen in #8822
• ci: run tests/cpu in a CPU-only lane on every PR Test run by @dierksen in #9293
• feat(cli): consistent kernel-wheel install options; docs: recommend download-kernels and document upgrades by @dierksen in #9328
• ci: catch provider CUDA-target mismatches on PRs (spec lint + cuobjdump check) by @dierksen in #8718
• Replace CLAUDE.md guidance with AGENTS.md by @saltyminty in #5328
• [SDSA] Let group_union_masks take more than 2^20 keys; move token_sparse into flashinfer.sdsa by @netanel-haber in #9323
• fix(gemm): add a cluster exit barrier to the 2-CTA TGV kernel by @zcnrex in #9146
• feat(moe): add SM12x NVFP4 fused MoE by @CarstyYou in #6510
• feat(moe_ep)!: select the NVLink one-sided sync protocol with sync_protocol, without CFT counted writes; name the flag-free protocol "sentinel" by @bobboli in #8770
• docs: reorganize AGENTS.md for clearer agent guidance by @righthandabacus in #9332
• docs: fix API documentation and improve checker guidance by @kangbintNV in #9307
• fix(attention): reuse paged cuDNN preparation on 9.24 by @YangXu1990uiuc in #9351
• feat(batch_invariant): batch-invariant bf16 MoE backend for the unified MoELayer (BatchInvariantBf16Config) by @bkryu in #9283
• docs(moe): regenerate the unified MoE activation matrix for the batch-invariant bf16 backend by @bkryu in #9369
• gdn: persist delta-rule DSL kernels in the CuTe-DSL disk cache by @aijanai in #9357
• feat(gemm): SM107 cute-dsl bmm_mxfp8 with MXFP8 output and strided A/out by @zcnrex in #9003
• fix(cake_muse_spark_norm): gated residual RMSNorm PDL programs read branch after the grid dependency wait by @yyihuang in #9359
• fix(moe): give the DSv3 routing kernel's PDL wait a memory clobber by @feih-nv in #9341
• fix(moe): initialise the DSv3 routing kernel's unused intermediate top-k slots by @feih-nv in #9344
• fix(trtllm-gen moe): publish routing outputs before the PDL trigger by @yechank-nvidia in #8610
• fix(moe): make zero tokens a no-op in the DSv3 routing launcher and check the launch by @feih-nv in #9345
• perf(moe_ep): faster SM90 FP8 pull MegaMoE autotuning (batched/rotated dispatch pull, blockwise N72/N80 tiles) by @ruoqianguo in #9209
• fix(attention): reject unsupported independent query lengths in paged fallback by @YangXu1990uiuc in #9371
• feat(attention): widen measured SM100 D128 paged auto domains by @YangXu1990uiuc in #9382
• refactor(attention): drop obsolete cuDNN Frontend compatibility paths by @YangXu1990uiuc in #9376
• perf(cake_muse_spark_norm): re-cut the gated residual RMSNorm PDL routes at 1-128 rows by @yyihuang in #9385
• feat(attention): admit measured SM100 D256 paged GQA by @YangXu1990uiuc in #9386
• fix(moe): check build before the pack contract in BatchInvariantMoeRunner by @Aneureka in #9391
• fix(prefill): validate FA2 prefill head dimensions by @lazypool in #4312
• Add FP8 per-(token, head) slot-mapping paged KV quantize by @ir1ka in #6863
• benchmarks: --backends batch_invariant in flashinfer_benchmark.py for the batch-invariant GEMM, quantization and sampling by @bkryu in #9368
• feat(moe_ep): order and deduplicate Rubin tuning candidates by @akaashrp in #9402
• fix(autotuner): keep the op's own bucket mapper under autotune(tuning_buckets=...) by @ji24077 in #8545
• fix(cake): ship build.json and csrc for the Muse Spark norm and TP MoE finalize backends by @zw-g in #9392

New Contributors

• @adsridhar made their first contribution in #6481
• @burak-inherent made their first contribution in #6132
• @michaelfeil made their first contribution in #8599
• @shyeh25 made their first contribution in #8744
• @tlysanhuo made their first contribution in #9148
• @smoothsmooth made their first contribution in #9247
• @SuperMarioYL made their first contribution in #4349
• @aijanai made their first contribution in #9357
• @yechank-nvidia made their first contribution in #8610
• @lazypool made their first contribution in #4312
• @zw-g made their first contribution in #9392

Full Changelog: v0.7.4rc2...v0.7.5rc1