flashinfer-ai/flashinfer

Summary top_k_sampling_from_probs and min_p_sampling_from_probs validate their per-request parameter tensor against the input batch but index it by the output row. When indices is longer than probs.size(0), every output row at or past probs.size(0) reads past the end of the parameter tensor and silently samples with a garbage top_k / min_p. check_tensor_param(maybe_top_k_arr, probs) — csrc/sampling.cu:216 — requires top_k_arr.size(0) == probs.size(0), via csrc/sampling_utils.h:34. Same for min_p at csrc/sampling.cu:262. The kernels read top_k_arr[bx] and min_p_arr[bx], where bx is blockIdx.x, i.e. the output row. The sibling kernels get this right and index by the input row: top_p_arr[row_idx] in TopPSamplingFromProbKernel, and top_k_arr[row_idx] / top_p_arr[row_idx] in TopKTopPSamplingFromProbKernel. So only the two [bx] sites are wrong. Reproduction On a B200 (SM100), 4 input rows, 16 output rows via indices = arange(16) % 4, vocab 512. The parameter tensor is a length-4 view c

Summary The Blackwell sm12x fused-MoE path compiles all of its CuTe-DSL kernels through explicit cute.compile(...), which unconditionally disables the DSL's file cache. The only caching layer is three in-process Python dicts, so every engine start re-pays the full MLIR JIT cost for every kernel specialization it touches — measured at ~40 s for the dynamic kernel's first hit and ~15–18 s per small-m micro-kernel specialization. In a serving context (vLLM --moe-backend=flashinfer_b12x) these compiles land mid-request and freeze the whole engine (vllm#47458: ~17.5 s stop-the-world stalls, all in-flight requests frozen including mid-decode). CUTE_DSL_CACHE_DIR / CUTE_DSL_DISABLE_FILE_CACHING have no effect on this path — the gate that would consult them is short-circuited before they are read. Environment flashinfer 0.6.12, 0.6.13, and 0.6.15.post1 (all verified; moe_dispatch.py cache structure unchanged across them) nvidia-cutlass-dsl 4.5.0 and 4.6.0 (both verified) SM120: RTX 5070

Noticed by @bkryu : "In our recent CI telemetry, we have been seeing that tests/gdn/test_decode_delta_rule.py runs very long (30-50 mins on every arch) and has been straining our CI. TL;DR of the cause: the reference and kernels are all fine. Most time is spent on cold kernel compilation. On a B200 I locally see ~800 cases that collapses to ~200 cute.compile calls based on the cache key logic. For some kernels, B/tile_v/T/cache flags/pool-strides each key a new compile. Would be great if we can make the kernels batch and/or seqlen dynamic to systematically reduce unit testing time."

FlashInfer — DGX Spark (SM121) Current Support Audit Date: 2026-04-24 Scope: Audit of SM121 (compute capability 12.1, DGX Spark) support across FlashInfer ops, APIs, backends, and testing. Goal: Improve Spark usability on FlashInfer by identifying existing gaps. Note: skip to "15. Consolidated Action Items" for TLDR. Architecture Context Gencode Target CUDA Requirement compute_120a / sm_120a RTX 5090 (SM120) CUDA 12.8 compute_120f / sm_120f SM12x family (120 + 121) CUDA 12.9 compute_121a / sm_121a DGX Spark (SM121) CUDA 12.9 SM120 vs SM121: Implications for FlashInfer Per the CUDA 13.2 Programming Guide — Compute Capabilities (Tables 29–33), SM120 and SM121 share a single "12.x" spec column: identical shared memory (99 KB/block, 100 KB/SM), warps (48/SM), registers (64K/SM), tensor core data types (TF32, BF16, FP16, FP8, FP6, FP4, INT8), TMA, and cluster support. Table 28 confirms compute_120f is compatible with both CC 12.0 and 12.1. The chip-level differences (SM

The doc string of trtllm_fp4_block_scale_moe indicates that the hidden_states can be bf16 and the hidden_states_scale is optional. But running the reproducing script hits error if passing hidden_states as bf16 and hidden_states_scale as None: File "/sgl-workspace/sglang/repro)fp4.py", line 303, in <module> main() File "/sgl-workspace/sglang/reprofp4.py", line 227, in main fi_outputs = trtllm_fp4_block_scale_moe( ^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/usr/local/lib/python3.12/dist-packages/flashinfer/fused_moe/core.py", line 2552, in trtllm_fp4_block_scale_moe return get_trtllm_moe_sm100_module().trtllm_fp4_block_scale_moe( ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/usr/local/lib/python3.12/dist-packages/flashinfer/fused_moe/core.py", line 1866, in trtllm_fp4_block_scale_moe_op intermediate_output = moe_op.trtllm_fp4_block_scale_moe( ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "python/tvm_ffi/c