Fix kernel launch config by murphymatt · Pull Request #5 · fw-ai/flashinfer

murphymatt · 2025-11-13T03:22:48Z

📌 Description

🔍 Related Issues

🚀 Pull Request Checklist

Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete.

✅ Pre-commit Checks

I have installed pre-commit by running pip install pre-commit (or used your preferred method).
I have installed the hooks with pre-commit install.
I have run the hooks manually with pre-commit run --all-files and fixed any reported issues.

If you are unsure about how to set up pre-commit, see the pre-commit documentation.

🧪 Tests

Tests have been added or updated as needed.
All tests are passing (unittest, etc.).

Reviewer Notes

## 📌 Description Brings in some changes to `test_hopper.py` to pass more unit tests * `test_deepseek_prefill` --> Raise tolerance for bf16 inputs * Others: The ``` token_pos_in_items_len=torch.tensor(token_pos_in_items_len) .to(dtype=torch.uint32) .to(0), ``` is an incorrect API and results in invalid input errors. Change it to: `token_pos_in_items_len=token_pos_in_items_len,` so that it matches the correct usage in e.g. [test_batch_prefill_kernels.py](https://github.com/flashinfer-ai/flashinfer/blob/6765cadd14fbedc9ffab428a87149a7d3f5d69f1/tests/attention/test_batch_prefill_kernels.py#L890) After this, `test_hopper.py` result improves to `3 failed, 2865 passed, 1320 skipped in 65.26s (0:01:05) `  ## 🔍 Related Issues  ## 🚀 Pull Request Checklist Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. ### ✅ Pre-commit Checks - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). ## 🧪 Tests - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.). ## Reviewer Notes

## 📌 Description Thor and Spark support when wheels are generating ## 🔍 Related Issues Output says that is not compatible. Only with JIT is working.  ## Summary by CodeRabbit * **New Features** * Broadened GPU architecture support to include additional newer architectures. * **Documentation** * Updated README and installation docs to show the revised CUDA architecture example list. * **Chores** * Adjusted release/nightly workflows and build scripts to select architectures using an expanded CUDA-version threshold and branching logic. * **Performance** * Extended architecture-specific build/runtime handling to cover an additional GPU architecture affecting memory-related behavior.  --------- Co-authored-by: Zihao Ye <expye@outlook.com> Co-authored-by: yzh119 <zihaoy@nvidia.com>

## 📌 Description Deprecate `tile_token_dim` in trtllm_moe. It is already not used and mark with deprecation warning, plan to deprecate totally in next major release  ## 🔍 Related Issues  ## 🚀 Pull Request Checklist Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. ### ✅ Pre-commit Checks - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). ## 🧪 Tests - [ ] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.). ## Reviewer Notes   ## Summary by CodeRabbit * **Refactor** * Removed the deprecated `tile_tokens_dim` parameter from MOE benchmarks and kernel functions, streamlining API calls and eliminating associated deprecation warnings.  Signed-off-by: jiahanc <173873397+jiahanc@users.noreply.github.com>

@Edenzzzz

Co-authored-by: @Edenzzzz ## 📌 Description Fixes flashinfer-ai/flashinfer#1022. Unlike flashinfer-ai/flashinfer#1231, this splits the inputs into separate prefill and decode inputs. It probably should be possible to automatically handle this splitting in Python so you can simply just provide a single batch of requests? To run the benchmark for this run: `python benchmarks/bench_mixed_attention.py` Performance: ===== Benchmark 1: (kv_len, qo_len) set ===== Prefill = 2 requests, 2048 Q len, 2048 KV len Decode = 128 requests, 2048 KV len Elapsed time (Batched Prefill): 0.65 ms Elapsed time (Batched POD Attention): 0.46 ms Elapsed time (Persistent BatchAttention): 0.56 ms **Batch POD speedup over Persistent BatchAttention: 1.22x** ===== Benchmark 2: (kv_len, qo_len) set ===== Prefill = 1 request, 2048 Q len, 2048 KV len Decode = 128 requests, 2048 KV len Elapsed time (Batched Prefill): 0.55 ms Elapsed time (Batched POD Attention): 0.41 ms Elapsed time (POD Attention): 0.41 ms Elapsed time (Sequential two kernels): 0.51 ms Elapsed time (Persistent BatchAttention): 0.45 ms **Batch POD speedup over Persistent BatchAttention: 1.11x** ===== Benchmark 3: (kv_len, qo_len) set ===== Prefill = 1 request, 4096 Q len, 4096 KV len Decode = 128 requests, 4096 KV len Elapsed time (Batched Prefill): 1.27 ms Elapsed time (Batched POD Attention): 0.86 ms Elapsed time (POD Attention): 0.82 ms Elapsed time (Sequential two kernels): 1.15 ms Elapsed time (Persistent BatchAttention): 1.08 ms Batch POD speedup over Persistent BatchAttention: 1.26x ===== Benchmark 4: (kv_len, qo_len) set ===== Prefill = 1 request, 4096 Q len, 4096 KV len Decode = 128 requests, 8192 KV len Elapsed time (Batched Prefill): 2.15 ms Elapsed time (Batched POD Attention): 1.52 ms Elapsed time (POD Attention): 1.54 ms Elapsed time (Sequential two kernels): 1.82 ms Elapsed time (Persistent BatchAttention): 1.76 ms **Batch POD speedup over Persistent BatchAttention: 1.16x** ===== Benchmark 5: (kv_len, qo_len) set ===== Prefill = 1 request, 6000 Q len, 7000 KV len Decode = 128 requests, 8192 KV len Elapsed time (Batched Prefill): 2.86 ms Elapsed time (Batched POD Attention): 2.03 ms Elapsed time (POD Attention): 1.95 ms Elapsed time (Sequential two kernels): 2.52 ms Elapsed time (Persistent BatchAttention): 2.45 ms **Batch POD speedup over Persistent BatchAttention: 1.20x** ## 🔍 Related Issues  ## 🚀 Pull Request Checklist Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. ### ✅ Pre-commit Checks - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). ## 🧪 Tests - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.). ## Reviewer Notes   ## Summary by CodeRabbit * **New Features** * Added a batched prefill+decode attention path with a public batch-oriented POD wrapper and JIT module export. * **Performance** * Benchmarks extended to include batched-path timings, memory bandwidth, elapsed-time and comparative speedup metrics across expanded prefill/decode scenarios. * **API** * Runtime binding for batched KV‑cache execution added; planning APIs now accept an optional colocated-CTA parameter that influences scheduling.  --------- Co-authored-by: Aditya K Kamath <akamath1997@gmail.com> Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com> Co-authored-by: Edenzzzz <wtan45@wisc.edu>

## 📌 Description Patch sm103 for 3xfp4 moe generation ## 🔍 Related Issues Following up of #2020 #1925 ## 🚀 Pull Request Checklist Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. ### ✅ Pre-commit Checks - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). ## 🧪 Tests - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.). ## Reviewer Notes ``` $ ls csrc/nv_internal/tensorrt_llm/cutlass_instantiations/103/gemm_grouped 100 103 80 $ pytest tests/moe/test_trtllm_cutlass_fused_moe.py 22 passed, 3 skipped, 1 warning in 771.89s (0:12:51) ```  ## Summary by CodeRabbit * **New Features** * Added support for Blackwell (SM103) GPU architecture in MOE (Mixture of Experts) operations with specialized CUTLASS-optimized modules.

## 📌 Description This PR does two things: * Add a check for the number of tokens and raise an exception if the max token size was exceeded * Adds an optional parameter to allow users to dial in an arbitrary workspace ## 🔍 Related Issues  ## 🚀 Pull Request Checklist Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. ### ✅ Pre-commit Checks - [ ] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [ ] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). ## 🧪 Tests - [x] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.). ## Reviewer Notes   ## Summary by CodeRabbit * **New Features** * Added an optional configurable workspace buffer size for all-reduce operations with a sensible default to preserve backwards compatibility. * Runtime input validation now enforces 2D inputs and token-count limits, with clearer error messages guiding corrective actions. * **Tests** * Expanded test coverage for workspace behavior: default sizing, explicit sizing, and negative tests for insufficient workspace. * Tests now allow supplying an explicit workspace size to validate allocation and reuse scenarios.

## 📌 Description - Small optimization for TRT-LLM Gen MoE finalize kernel TopK=8, NumExperts=128, HiddenSize=4096 | BS | Baseline, us | Optimized, us | Speed-up | | ------------- | ------------- | ------------- | ------------- | | 256 | 11 | 6 | 1.83 | | 512 | 12 | 7 | 1.71 | | 1024 | 16 | 15 | 1.06 | | 4096 | 55 | 49 | 1.12 | | 8192 | 107 | 95 | 1.13 | | 16384 | 205 | 183 | 1.12 |  ## 🔍 Related Issues  ## 🚀 Pull Request Checklist Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. ### ✅ Pre-commit Checks - [ ] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [ ] I have installed the hooks with `pre-commit install`. - [ ] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). ## 🧪 Tests - [ ] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.). ## Reviewer Notes   ## Summary by CodeRabbit * **New Features** * Enabled vectorized, Top-K unrolled finalize path for MOE (Mixture of Experts) kernel operations with improved performance. * Added support for multiple data types (bfloat16, float, half) with enhanced type specialization and packing. * Introduced runtime validation for TopK configurations (≤ 64) to ensure optimal vectorized execution.

## 📌 Description Refactor fused_moe test. Split test on model+precision. Part [1]: - test deepseek (kimi, lite) fp8 block-scaled fused moe - default TP8 - PDL enabled - MajorK weight layout - higher tolerance and matching percentage Next Part [2]: - add BlockMajorK weight layout Next Part [x]: - Per Tensor FP8 MoE, FP4MoE later: - refactor llama4, topk?, renormalize? routing tests ## 🔍 Related Issues  ## 🚀 Pull Request Checklist Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. ### ✅ Pre-commit Checks - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). ## 🧪 Tests - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.). ## Reviewer Notes   ## Summary by CodeRabbit * **Tests** * Added a comprehensive FP8 block-scale fused Mixture-of-Experts test validating end-to-end correctness across many routing, expert and precision configurations. Includes randomized inputs, per-token/per-expert workflows, extensive parameterizations, diagnostic statistics, autotune-path checks, and a minimal sanity run.

* chore: rename FLASHINFER_JIT_VERBOSE to FLASHINFER_JIT_DEBUG for clarity (#1946)  Rename environment variable `FLASHINFER_JIT_VERBOSE` to `FLASHINFER_JIT_DEBUG` to better reflect its actual behavior. - `FLASHINFER_JIT_DEBUG`: Enable debug mode during compilation (disable optimization, add debug symbols) - The previous name `FLASHINFER_JIT_VERBOSE` implied "showing more compilation info", which was confusing - Maintained backward compatibility: falls back to `FLASHINFER_JIT_VERBOSE` if `FLASHINFER_JIT_DEBUG` is not set  Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [x] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.).   * **Refactor** * Introduced FLASHINFER_JIT_DEBUG environment variable for controlling JIT debug builds with backward compatibility for legacy FLASHINFER_JIT_VERBOSE. * Enhanced debug build configuration with improved compiler and CUDA debugging flags. Non-debug builds continue using -O3 optimizations.  * fix: Fix trtllm-gen prefill IMA when batch_size==1 (#1912)   Current PR fixes the test and benchmark codes IMAs when running trtllm-gen paged & ragged prefill with batch size 1 -- the issue was described in https://github.com/flashinfer-ai/flashinfer/issues/1898 Root cause of the issue: `flashinfer.prefill.trtllm_ragged_attention_deepseek` and `flashinfer.prefill.trtllm_batch_context_with_kv_cache` both require `max_q_len` to match the length of the query when batch size is 1. **Updated PR:** Issue has been addressed from the kernel-side so that the "*`max_q_len` to match the length of the query when batch size is 1*" is no longer required. Current PR updates trtllm-gen FMHA cubins to latest and brings minor updates to kernel metadata. Unit test results after PR: ``` $ pytest tests/attention/test_trtllm_gen_attention.py ... platform linux -- Python 3.12.11, pytest-8.4.2, pluggy-1.6.0 rootdir: /flashinfer configfile: pytest.ini collected 2320 items ... 2055 passed, 264 skipped, 1 xfailed in 224.43s (0:03:44) ``` **Description of previous solution:** ~~Updating `max_q_len` to `cum_seq_lens_q[-1].item()` within the `trtllm_ragged_attention_deepseek` or `trtllm_batch_context_with_kv_cache` functions are not a viable option because the CPU-side synchronization breaks the deterministic and fully device-side execution required during CUDA graph capture. The workaround was thus to update the test & benchmark codes that call the trtllm prefill functions, and clearly state in the docstring that when batch_size == 1, max_q_len must match the query size.~~ https://github.com/flashinfer-ai/flashinfer/issues/1898  Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [ ] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.).   * **Bug Fixes** * Removed the automatic batch_size=1 restriction for a native backend, enabling its use in more scenarios while other constraints remain. * **New Features** * Added configurable block-sparse attention support to kernel parameters. * **Documentation** * Clarified supported attention optimizations and backend capabilities in the benchmarks docs. * **Tests** * Expanded tests with configurable sequence lengths and added dedicated batch-size-1 test coverage.  --------- Co-authored-by: Zihao Ye <expye@outlook.com> * Feature: Support Relu2 activation in fused MoE (#1954) Added support for Relu2 activation in cutlass fp8 FusedMoE path. `Relu2(x) = Relu(x)^2`. Validated this works correctly on H100 and B200.  * **New Features** * Added Relu2 as a selectable activation across MOE operations and exposed activation_type configuration to public MOE APIs and runner interfaces (Swiglu remains the default). * **Behavior** * Certain GEMM execution paths now explicitly reject Relu2 and raise a clear runtime error instead of silently failing.  --------- Signed-off-by: Amir Klein <203507526+amirkl94@users.noreply.github.com> * fix: Add cutlass as an mm_fp4 backend in compute capability 12.0 in benchmark code (#1959)  Previously `backend='cutlass'` was not available to be benchmarked in `flashinfer_benchmark.py` for compute capability 12.0 while the kernel actually has been available. Current PR marks the backend as available. Example output of being runnable after PR: ``` [INFO] args = Namespace(routine='mm_fp4', no_cuda_graph=False, use_cupti=False, refcheck=True, allow_output_mismatch=False, random_seed=42, verbose=2, output_path=None, num_iters=30, dry_run_iters=5, case_tag=None, generate_repro_command=False, repro_command='', batch_size=1, m=1024, n=7168, k=512, tile_size=128, group_size=1, scale_major_mode='MN', input_dtype='fp8_e4m3', mat2_dtype='fp8_e4m3', out_dtype='bfloat16', mma_sm=1, backends=['cudnn', 'cutlass', 'trtllm'], use_128x4_sf_layout=True, use_nvfp4=True, autotune=False) [INFO] Running testMmFp4 [INFO] FlashInfer version: 0.4.1 [VVERBOSE] gpu_name = 'NVIDIA_RTX_PRO_6000_Blackwell_Server_Edition' [WARNING] trtllm for routine mm_fp4 is not supported on compute capability 12.0. Skipping. [VVERBOSE] input_fp4.shape = torch.Size([1024, 256]) [VVERBOSE] input_fp4.dtype = torch.uint8 [VVERBOSE] mat2_fp4.shape = torch.Size([7168, 256]) [VVERBOSE] mat2_fp4.dtype = torch.uint8 [PERF] cudnn :: median time 0.014 ms; std 0.000 ms; achieved tflops 535.891 TFLOPs/sec; achieved tb_per_sec 1.196 TB/sec [PERF] cutlass :: median time 0.015 ms; std 0.000 ms; achieved tflops 515.203 TFLOPs/sec; achieved tb_per_sec 1.150 TB/sec ```   Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.).   * **Chores** * Expanded backend support for benchmarking routines on compute capability 12.0, adding compatibility with additional processing backends.  * rebase on fw repo branch * unittest: fix deepgemm sha256 (#1953)  Deepgemm unittest failed because of out-dated sha256, this PR fixes the issue.  Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [ ] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [ ] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.).   * **Chores** * Updated internal artifact version information to support latest optimizations and improvements.  * misc: Update artifacts docstring and MetaInfoHash (#1967)  Amendment to [PR 1761](https://github.com/flashinfer-ai/flashinfer/pull/1761), appending docstring to two artifactory path classes and deprecating need to update MetaInfoHash by directly accessing the checksum.txt file.  Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [ ] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [ ] I have installed the hooks with `pre-commit install`. - [ ] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [ ] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.).   * **New Features** * Added runtime integrity checks for compiled artifacts that verify and use checksum data during loading to prevent missing or mismatched artifact headers. * **Refactor** * Switched artifact hash resolution to compute hashes dynamically from provided checksums, improving validation, reliability, and resilience when loading precompiled components.  * silu_and_mul nvfp4 quanization fusion rework (#1927)  This PR reverts https://github.com/flashinfer-ai/flashinfer/pull/1774 and https://github.com/flashinfer-ai/flashinfer/pull/1835 which have some issues with some shapes under cuda graph. The kernels ported in this PR comes from SGLANG. [[NVIDA] [1/N] Nvfp4 Masked Gemm: Add quant op for the flashinfer grouped gemm](https://github.com/sgl-project/sglang/pull/9200/files) and [[NVIDIA] [2/N] Optimize silu_and_mul_scaled_fp4_grouped_quant perf](https://github.com/sgl-project/sglang/pull/9556/files) by @kaixih .  Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [ ] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [ ] I have installed the hooks with `pre-commit install`. - [ ] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [ ] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.).   * **New Features** - Added grouped FP4 quantization (scaled_fp4_grouped_quantize) and an NV-focused Silu+Mul expert quantization entry (silu_and_mul_scaled_nvfp4_experts_quantize). * **API Changes** - Replaced legacy batched APIs with new expert/grouped APIs; removed legacy mask parameter from FP4/MXFP8 quantization signatures and adjusted FP4 output layouts/types. * **Documentation** - Updated docs to list new functions and remove deprecated symbols. * **Tests** - Updated tests to validate new quantization paths, shapes, dtypes, and layouts.  --------- Signed-off-by: Shu Wang. <shuw@nvidia.com> * unittest: fix test_artifacts.py (#1950) * chore: update the list of authorized codeowners (#1970)  Add @djmmoss @jiahanc to the authorized codeowner list.  Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [ ] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [ ] I have installed the hooks with `pre-commit install`. - [ ] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [ ] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.).   * **Chores** * Updated internal codeowner authorization configuration.  * Added heuristic for trtllm_allreduce_fusion (#1972)  The original heuristic does not accurately reflect the performance of oneshot/twoshot. Updated with heuristics based on this benchmark [allreduce_test.py](https://github.com/user-attachments/files/23094671/allreduce_test.py). The benchmark uses hidden_dim of Llama3, LLama4 and GPT-OSS and combinations of token_num, fusion patterns and fp32_acc. The results are at the bottom. TL;DR token_num is a bad predictor of whether to use oneshot or twoshot. Using the communication size of oneshot is a good predictor, but only if we treat each TP separately. Fusion patterns and fp32_acc is irrelevant to the choice. <img width="1800" height="3600" alt="comm_size_TP=2" src="https://github.com/user-attachments/assets/2874157e-6268-421a-8f45-00491b652702" /> <img width="1800" height="3600" alt="comm_size_TP=4" src="https://github.com/user-attachments/assets/2cdfdb9d-569e-401b-89ad-787f8d755ac1" /> <img width="1800" height="3600" alt="comm_size_TP=8" src="https://github.com/user-attachments/assets/fbb147da-3479-4dbc-85b8-c27a735d0cd6" /> <img width="1800" height="3600" alt="comm_size_Enlarge_TP=2" src="https://github.com/user-attachments/assets/e070c81f-edf9-4d7f-ab95-fa6dea9f42f2" /> <img width="1800" height="3600" alt="comm_size_Enlarge_TP=4" src="https://github.com/user-attachments/assets/3b1c51d2-56ca-4d34-9bfd-8082390cc95e" /> <img width="1800" height="3600" alt="comm_size_Enlarge_TP=8" src="https://github.com/user-attachments/assets/9a8095b4-11bc-4021-80c6-f2be69b33021" /> <img width="1800" height="3600" alt="comm_size_TP=248" src="https://github.com/user-attachments/assets/66956ebe-6cf0-43e8-93ce-950b1079148a" /> <img width="1800" height="3600" alt="comm_size_Enlarge_TP=248" src="https://github.com/user-attachments/assets/0cd6982c-da42-4f42-b0ad-5ef564b2e78e" /> <img width="1800" height="3600" alt="token_num_TP=248" src="https://github.com/user-attachments/assets/2968ca7c-2059-4305-8e4d-5b70a32faaee" /> <img width="1800" height="3600" alt="token_num_Enlarge_TP=248" src="https://github.com/user-attachments/assets/881ba86d-fc71-4cbc-b5a6-c050f255d618" />   Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [ ] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [ ] I have installed the hooks with `pre-commit install`. - [ ] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [ ] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.). --------- Co-authored-by: yzh119 <zihaoy@nvidia.com> * Bump tvm ffi to stable version 0.1.0 (#1960)  This PR bumps the tvm-ffi to stable version 0.1.0 and update the flashinfer code base.  https://github.com/flashinfer-ai/flashinfer/pull/1939 Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [ ] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [ ] I have installed the hooks with `pre-commit install`. - [ ] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [ ] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.).   * **Chores** * Relaxed build dependency pins for apache-tvm-ffi and setuptools across project configs; removed installation of multiple build packages from the nightly CI step. * **Refactor** * Modernized internal CUDA/tensor access patterns to a consistent accessor API across many modules. * **Bug Fixes** * GEMM runner now returns the output tensor in the correct (non‑transposed) orientation.  --------- Co-authored-by: Zihao Ye <expye@outlook.com> Co-authored-by: yzh119 <zihaoy@nvidia.com> * Update Docker CI tags to 20251024-0e48aaf (#1975) This PR updates the Docker CI image tags to the latest version: `20251024-0e48aaf` Updated images: - flashinfer/flashinfer-ci-cu126:20251024-0e48aaf - flashinfer/flashinfer-ci-cu128:20251024-0e48aaf - flashinfer/flashinfer-ci-cu129:20251024-0e48aaf - flashinfer/flashinfer-ci-cu130:20251024-0e48aaf Auto-generated by [release-ci-docker workflow](https://github.com/flashinfer-ai/flashinfer/actions/runs/18778064727)  * **Chores** * Updated CI/CD Docker image configurations to ensure consistency and reliability across build environments.  Co-authored-by: yzh119 <11773619+yzh119@users.noreply.github.com> * fix: Make attention microbenchmark correctly use page table (#1976)  Current microbenchmark code does not provides instantiated `block_tables` to all backends. The omission had no impact to correctness or perf because page tables are instantiated linearly when not provided, but will manifest as mismatches if it is shuffled. The current PR simply calls the FlashInfer APIs in their intended way. **No changes to library code**   Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.).   * **Refactor** * Enhanced consistency in attention computation by aligning page-table parameter handling across different inference backend implementations for improved paged key-value cache operations.  * fix: Skipping attention sink Blackwell test outside of Blackwell (#1978)  `test_attention_sink_blackwell.py` checks `flashinfer.prefill.trtllm_batch_context_with_kv_cache` and `flashinfer.decode.trtllm_batch_decode_with_kv_cache` which are only supported on Blackwell SM100 and SM103. Existing check only skips testing of SM 11x or 12x, which causes failures on Hopper SM90. Test outputs: * H200: * Before Fix: `144 failed, 1 warning in 9.20s` * After Fix: `144 skipped, 1 warning in 0.42s` * B200: * After Fix: `144 passed in 34.64s `   Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.).   * **Tests** * Updated GPU compatibility checks for attention sink tests to target specific GPU architectures (SM100/SM103). Tests now run exclusively on supported GPU models with updated filtering criteria.  * feat: enable deepgemm jit for fp8 block-scale on SM90 (#1969)  Enable JIT compile for the FP8 DeepGEMM kernels, NVRTC is currently disabled it uses NVCC by default. Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.).  * **Refactor** * JIT include directory discovery now uses the flashinfer-python package instead of the previous package. * Updated resolved include path to the flashinfer data location. * Runtime compilation now consistently uses NVCC; the prior environment-variable toggle was removed. * Updated warning text when the expected package installation cannot be found.  --------- Signed-off-by: Duncan Moss <djm.moss@gmail.com> * chore: Update CODEOWNERS (#1949) This PR updates the CODEOWNERS file based on git commit history analysis from the last 180 days. - Updated `.github/CODEOWNERS` with current code ownership based on: - Commit frequency - File coverage - Commit recency 1. Review the changes to `.github/CODEOWNERS` 2. Verify that the assigned owners are appropriate for each module 3. Make manual adjustments if needed before merging - This is an automated PR generated weekly - Minimum commits threshold: 1 - Analysis period: 180 days - Directory depth: 3 levels - Top N owners per module: 5 --- 🤖 This PR was automatically generated by the [update-codeowners workflow](.github/workflows/update-codeowners.yml)  * **Chores** * Updated internal code ownership assignments. --- **Note:** This update contains no user-facing changes or feature updates. It is an internal administrative modification.  Co-authored-by: flashinfer-bot <flashinfer-bot@users.noreply.github.com> Co-authored-by: Claude <noreply@anthropic.com> * fix: correct PDL parameter handling in RopeQuantize kernel (#1982)  - **Issue**: The `stream` parameter was being passed to the wrong position in the `RopeQuantize` function call due to missing `enable_pdl` parameter. SGLang will hang before this pr. - **Fix**: Added the `enable_pdl` parameter to the function signature and properly aligned all parameters - **Issue**: When `enable_pdl=true`, the kernel would throw CUDA errors due to incorrect PDL attribute handling - **Fix**: Aligned the implementation with `csrc/fmhaReduction.cu`.   Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.).   * **New Features** * Added PDL (Programmatic Dynamic Launch) benchmarking capability for rope quantization operations. * Extended configuration options to enable or disable PDL functionality. * **Tests** * Updated test suite to validate PDL enabled and disabled scenarios in rope quantization workflows.  * Fix: Verify scales are not None for Cutlass FP8 FusedMoE (#1961) Verify quant scales for fp8 are non null in cutlass FusedMoE path. Currently, if these tensors are passed as None from python it will result in segmentation fault.  * **Bug Fixes** * Enhanced validation for FP8 quantization parameters to improve system robustness and prevent potential null reference issues during quantization operations, reducing the risk of runtime errors when processing quantized model data.  --------- Signed-off-by: Amir Klein <203507526+amirkl94@users.noreply.github.com> * add xqa fp8 mha and fp8 kv cache (#1769)  Add xqa fp8 mha and fp8 kv cache. Add fp8 mla for sm120. Use vllm kv layout.  Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.).   * **New Features** * MLA-based attention path and dedicated MLA entrypoints (SM120/121) * FP8 KV-cache support with optional paged KV layout and separate K/V cache inputs * Asynchronous tensor-map/TMA and matrix-descriptor primitives for high-throughput GPU transfers * Dtype-driven config and expanded GPU SM gating for builds/runtimes * **Bug Fixes** * Improved numerical stability for attention mask initialization * **Tests** * Expanded coverage for MLA, FP8, FP16/BF16, and new cache layouts * **Documentation** * Added XQA API docs and new public symbols  --------- Signed-off-by: Qidi Sang <200703406+qsang-nv@users.noreply.github.com> Co-authored-by: yzh119 <zihaoy@nvidia.com> * unittest: fix failed unittest on hopper (#1952)  Some invalid configuration are generated in JIT warmup (mixed precision) function `gen_prefill_attention_modules`.  Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [x] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.).   * **Tests** * Updated test infrastructure to enhance compatibility handling for specific hardware acceleration scenarios, improving test robustness for mixed-precision configurations.  * docs: Update documented versioning scheme to right-shifted semver (#1990)  Based on discussion with @yzh119 and others, we're planning to follow the vLLM "right-shifted" versioning scheme. This PR updates the docs to reflect that. Previously we said we would follow Semantic Versioning (see #1553). However, we recently re-considered this approach, to better match the conventions followed by vLLM and PyTorch. Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). Docs only, so no new tests are needed. Did not verify passing unit tests. - [x] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.).   * **Documentation** * Updated release versioning scheme to a "right-shifted" format (major.minor.patch[.post1]) with an optional post-release suffix for expedited follow-up releases.  * Bugfix: Change get() -> GetDLTensorPtr() in cutlass FusedMoE validations (#1995) Using different API after `apach-tvm-ffi` version bump.  * **Bug Fixes** * Improved null-pointer validation for FP8 quantization tensors used during inference, increasing robustness and reducing risk of runtime errors.  --------- Signed-off-by: Amir Klein <203507526+amirkl94@users.noreply.github.com> * unittest: Add SM arch checks to skip unsupported tests on Hopper (#1998)  A number of unit tests fail on Hopper because they either do not have a support-check or fail based on "what is not supported" while missing SM90. Current PR adds checks based on "what is supported" and skips if not in the supported list of SMs. Special case of `mm_fp4` where `mm_fp4.is_backend_supported(backend, compute_capability_number)` now exists and is used to skip tests if not supported. Impacted tests: * tests/attention/test_trtllm_gen_attention.py * tests/attention/test_trtllm_gen_mla.py * tests/gemm/test_bmm_fp8.py * tests/gemm/test_mm_fp4.py * tests/gemm/test_groupwise_scaled_gemm_fp8.py * tests/gemm/test_groupwise_scaled_gemm_mxfp4.py * tests/moe/test_trtllm_gen_fused_moe.py   Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.).  * Added workspace check and reflected this in test (#1991)  This PR attempts to fix #1986 (to be confirmed by requester) The issue is that num_tokens was larger than MAX_TOKEN_NUM, which results in an IMA, or even in a hang. To address this, I added a validation check. This required a non-breaking API change: * create_ipc_workspace_for_all_reduce_fusion now has an optional "create_metadata" bool, which results in an additional return value * it is made optional as additional return value could break the API * trtllm_allreduce_fusion now takes an optional metadata dictionary * When provided, this will run the validation check * again, this is also optional, to avoid breaking the api In addition this PR deprecates the older AllReduce functionality so it can be removed in a major version bump.  Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [ ] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [ ] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [x] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.).   * **API Changes** * Workspace creation can optionally return metadata describing the workspace configuration (create_metadata flag). * Allreduce fusion operations accept optional metadata to validate runtime parameters against the workspace and raise clear errors on mismatch. * A workspace destruction endpoint was renamed for naming consistency. * Legacy wrappers were marked deprecated and now point users toward the newer fusion variants.  * minor fix for xqa (#1994)   1 change xqa_mla comments to be consistent with mla instead of mha. 2 put cudaMemcpyFromSymbol/cudaFuncSetAttribute outside of launch function to avoid breaking cuda graph capture 3 use int32 as pagetable index  Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [x] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.).   * **New Features** * Added MLA variant documentation clarifying SM120 GPU requirement and fixed head group ratio configuration. * **Documentation** * Updated data type specifications for XQA operations; page table now requires int32 instead of uint32. * Added max sequence length derivation notes for page-table-based configurations. * Clarified MLA variant input/output data types (float8_e4m3fn and bfloat16). * **Bug Fixes** * Corrected data type handling in page table processing to ensure compatibility.  Signed-off-by: Qidi Sang <200703406+qsang-nv@users.noreply.github.com> Co-authored-by: yzh119 <zihaoy@nvidia.com> * Feature: Add support for L40 FusedMoE in cutlass path (#1973) Fixed a few compilation issues for L40, and removed 1 gemm tactic for `sm == 89` that crashes due to: ``` Assertion failed: GPU lacks the shared memory resources to run GroupedGEMM kernel ``` Ran `pytest tests/moe/test_trtllm_cutlass_fused_moe.py` manually on an L40 GPU and verified all tests passed.  * **New Features** * Official support for SM89 target: build/JIT flags and a public generation path to target it. * **Bug Fixes / Compatibility** * Clarified FP8/FP4 dispatch: FP8 paths enabled for SM89; FP4 usage remains gated and now requires explicit enablement. * **Performance** * Adjusted kernel/tile selection order for certain FP8 paths to prefer SM89-optimized options. * **Chores** * Reduced logging severity for failed tactic profiling to warn/debug.  --------- Signed-off-by: Amir Klein <203507526+amirkl94@users.noreply.github.com> * unittest: Add head dim 256 test cases and mark as xfail (#1999) * feat: autotune tile_tokens_dim in trtllm-gen MOE (#1980)  - Update the autotune logic in trtllm-gen moe. Instead of using a fixed `tile_tokens_dim`, tune in a range of `[max(8,tile_token_dim/2), tile_token_dim, min(128, tile_token_dim*2), min(128, tile_token_dim*4)]` - Add FP8 MOE autotune logic, initial PR https://github.com/flashinfer-ai/flashinfer/pull/1494 from @aleozlx, update logic to sync with new autotuner. - Update logic in `test_trtllm_gen_fused_moe.py`. - Update the `conftest.py` to speed up test, previously use `try_first` which introduce duplicate run - Add log_once in logger   Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.).   * **New Features** * Runtime autotuning with per-tile dynamic routing and selectable MoE runner options (gated activation, shuffled-weight, weight-layout). * One-time (deduplicated) logging helpers added to JIT logger. * **Deprecations** * tile_tokens_dim removed from new paths and marked deprecated in legacy entry points; new tuning parameters introduced for autotuning. * **Tests** * Tests refactored for autotuning/routing with new helpers and improved handling/reporting for missing JIT cache.  --------- Signed-off-by: jiahanc <173873397+jiahanc@users.noreply.github.com> Co-authored-by: yzh119 <zihaoy@nvidia.com> * Fix trtllm-gen attention illegal memory access (#2002)  This PR fixes illegal memory access of trtllm-gen attention kernels. It changes the workspace buffer from `int_workspace_buffer` to `float_workspace_buffer`. `int_workspace_buffer` is a fixed sized buffer and not initialized to zero, which should not be used.  Issue #1928 Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.).   * **Bug Fixes** * Fixed memory allocation in the decode module to improve computation accuracy and stability during text generation.  * release: Bump version for v0.5.0rc1 release; (#2008)  Update version in `version.txt` to v0.5.0 as we prepare for v0.5.0rc1 release.   Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.).   * **Chores** * Version bump to 0.5.0 (no functional changes)  * bugfix: fix regex in update wheel index script (#2009)  The regex cannot recognize release candidates (`v0.5.0rc1`) or post releases (`v1.2.3.post1`): https://github.com/flashinfer-ai/flashinfer/actions/runs/18929490991/job/54049304551 This PR fixes the issue. https://github.com/flashinfer-ai/flashinfer/actions/runs/18929490991/job/54049304551 Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [x] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.).   * **Chores** * Enhanced version string parsing in the wheel package indexing process to support more complex version formats, including pre-release, post-release, and development versions, ensuring compatibility with PEP 440 versioning standards.  * fix: Enable SM121 for mm_fp4 (#2012)  In #1809 we previously added a compute-capability-based support check for `mm_fp4`. However, we missed enabling SM121 for backend = `cudnn` and `cutlass`. Additionally, we marked `trtllm` as supported on SM120 when it is not. Current PR fixes it. Example benchmark and pytest command on SM121 after the fix ``` (py312) root@f414f262f02a:/flashinfer/benchmarks# python3 flashinfer_benchmark.py --routine mm_fp4 --m 8192 --n 7168 --k 512 --out_dtype bfloat16 --backends cudnn cutlass --use_128x4_sf_layout --use_nvfp4 --refcheck --use_cupti /opt/conda/envs/py312/lib/python3.12/site-packages/torch/cuda/__init__.py:285: UserWarning: Found GPU0 NVIDIA GB10 which is of cuda capability 12.1. Minimum and Maximum cuda capability supported by this version of PyTorch is (8.0) - (12.0) warnings.warn( [PERF] cudnn :: median time 0.656 ms; std 0.025 ms; achieved tflops 91.701 TFLOPs/sec; achieved tb_per_sec 0.185 TB/sec [PERF] cutlass :: median time 0.669 ms; std 0.022 ms; achieved tflops 89.859 TFLOPs/sec; achieved tb_per_sec 0.181 TB/sec (py312) root@f414f262f02a:/flashinfer# pytest tests/gemm/test_mm_fp4.py ====================================================================================================================== test session starts ====================================================================================================================== platform linux -- Python 3.12.11, pytest-8.4.2, pluggy-1.6.0 rootdir: /flashinfer configfile: pytest.ini collected 3240 items ... ======================================================================================================================= warnings summary ======================================================================================================================== ../opt/conda/envs/py312/lib/python3.12/site-packages/torch/cuda/__init__.py:285 /opt/conda/envs/py312/lib/python3.12/site-packages/torch/cuda/__init__.py:285: UserWarning: Found GPU0 NVIDIA GB10 which is of cuda capability 12.1. Minimum and Maximum cuda capability supported by this version of PyTorch is (8.0) - (12.0) warnings.warn( -- Docs: https://docs.pytest.org/en/stable/how-to/capture-warnings.html ========================================================================================================= 450 passed, 2790 skipped, 1 warning in 8.24s ========================================================================================================== ```   Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [ ] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.).   * **New Features** * Expanded hardware compatibility by adding support for newer NVIDIA GPU architectures. * FP4 quantized operations now available across multiple backends on supported devices.  * fix: ensure SM120/121 SFA/SFB contiguity (#1963)  Fix the regression in vLLM and SGLang with FI 0.4.0 in bmm_fp8  Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [ ] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [ ] I have installed the hooks with `pre-commit install`. - [ ] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [ ] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.). cc: @yzh119  * **Bug Fixes** * Fixed memory layout handling for tensor operations in GPU computations to ensure proper alignment, improving stability and performance.  * More realistic bench for POD Attn (#2013)  Use real head sizes, seq lens and add comparison with sequential prefill + decode. Results on H100 (without overlap, which only adds ~150GB/s for persistent): <img width="433" height="571" alt="image" src="https://github.com/user-attachments/assets/50de01cd-e5ca-450c-9cc0-521d83b7e487" /> cc @yzh119  Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [ ] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [ ] I have installed the hooks with `pre-commit install`. - [ ] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [ ] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.).   * **New Features** * Added comprehensive performance benchmarking for batch attention operations with detailed timing measurements. * Introduced sequential dual-kernel benchmark path with extended memory bandwidth reporting. * **Tests** * Updated benchmark test configurations to use deterministic, fixed values for improved reproducibility. * Adjusted benchmark parameters for consistency across test iterations.  * Feature: Support non-gated activation in cutlass fused MoE nvfp4 (#2011) This PR removes an assertion in the cutlass fused moe bindings to enable non-gated activations in nvfp4. It also adds a test for this path with relu2 activation. N/A Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [v] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [v] I have installed the hooks with `pre-commit install`. - [v] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [v] Tests have been added or updated as needed. - [v] All tests are passing (`unittest`, etc.). N/A  * **New Features** * Enhanced quantized Mixture of Experts models to support configurable activation types (Swiglu and ReLU2) in the NVFP4 quantization path. * Improved parameter handling to correctly adapt weight shapes and quantization settings based on the selected activation type.  --------- Signed-off-by: Omer Ullman Argov <118735753+omera-nv@users.noreply.github.com> * feat: add xqa backend and completes NHD/HND coverage for trtllm-gen/xqa backend (#2001)  Expose xqa backend to trtllm attention interface, and improve layout coverage of trtllm-gen and xqa backends. Now both trtllm-gen/xqa supports NHD/HND kv-cache layout. * support NHD layout for trtllm-gen * refactor xqa (https://github.com/flashinfer-ai/flashinfer/commit/869c0c1c6bc199f82f30c23ab78a1b4aa9a1bd3a) * allow user passed stride_page/head/token * support both HND and NHD * remove macros such as PAGED_KV_CACHE_LAYOUT and USE_PAGED_KV_CACHE * adding unittests for both trtllm-gen/xqa on NHD/HND * adding unified API for trtllm-gen/xqa, and unified unittest  Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [x] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.).   * **New Features** * Added xqa-based batch decode API and public kv_layout option (NHD/HND); added enable_pdl toggle to inference wrappers. * **Improvements** * Automatic backend selection for decoding, consistent KV-layout normalization across paths, and unified stride-aware paged-KV handling with layout-aware shapes, scales, and workspace handling. * **Tests** * Expanded tests to cover both KV layouts, enable_pdl, new batch-decode workflows, backend/layout permutations, and fp8/mixed-dtype scenarios.  --------- Signed-off-by: Qidi Sang <200703406+qsang-nv@users.noreply.github.com> Co-authored-by: yzh119 <zihaoy@nvidia.com> Co-authored-by: Zihao Ye <expye@outlook.com> * test: Enable xfailed trtllm decode long seqlen tests and update microbenchmark (#2018)  [tests/attention/test_trtllm_gen_attention.py](https://github.com/flashinfer-ai/flashinfer/blob/v0.5.0rc2/tests/attention/test_trtllm_gen_attention.py#L1021-L1076) was failing and therefore marked xfail. PR #2002 fixed the underlying root cause. Current PR thus removed the `xfail` marker so that these long seqlen cases could be fixed moving forward. Additionally, PR #2002 revealed a bug in the microbenchmark script where [trtllm_batch_decode_with_kv_cache](https://github.com/flashinfer-ai/flashinfer/blob/v0.5.0rc2/flashinfer/decode.py#L2082-L2083) explicitly requires the workspace to be zeroed before first use: ``` workspace_buffer : torch.Tensor. Must be initialized to 0 for its first use. workspace ``` while the microbenchmark code does not zero out, causing undefined behavior such as IMAs that depend on the ordering of backends tested. Current PR fixes the issue by explicitly calling `workspace_buffer.zero_()` between testing different backends.   Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [ ] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.).   * **Bug Fixes** * Improved stability of performance benchmarks by properly resetting workspace buffer between backend invocations. * **Tests** * Enabled previously skipped test for long sequence length handling.  * Updated decorator to support unspecified default (#2026)  Updated decorator to support unspecified default. This was causing issues when calling mm_fp4 without backend specified. Also added SM 110 as a supported backend on the cutlass backend (mm_fp4)  Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [ ] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [ ] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [ ] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.).   * **New Features** * FP4 Cutlass GEMM now supports the SM110 GPU compute capability. * **Bug Fixes** * Kernels called without an explicit backend now consistently use the default backend. * **Tests** * Added a unit test to verify default backend selection and correct results when backend is omitted.  * release: Bump version for v0.5.1 release (#2031)  Update `version.txt`   Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.).   * **Chores** * Version updated to 0.5.1  * ci: Update cudnn version requirements in CI container (#2039)  cuDNN versions specified in CI container setup (`docker/install/install_python_packages.sh`) are currently 9.11 and 9.12. In unit testing, this causes issues as `mm_fp4(backend='cudnn')` is not supported on Spark (sm121) for older cuDNN versions in cu130. Failure is due to cuDNN version shipped with container being too old. In the [latest container build pipeline output](https://github.com/flashinfer-ai/flashinfer/actions/runs/18778064727/job/53577233568#step:6:727), cudnn 9.13.0.50 is installed ``` ``` Current PR updates the minimum cudnn version for both [cu12](https://pypi.org/project/nvidia-cudnn-cu12/#history) and [cu13](https://pypi.org/project/nvidia-cudnn-cu13/#history) to 9.14.0.64. cudnn 9.13 --> unit test fails with 180 failed, 270 passed, 2790 skipped, 1 warning in 8.97s ``` =================================================================================================================================================== test session starts =================================================================================================================================================== platform linux -- Python 3.12.11, pytest-8.4.2, pluggy-1.6.0 rootdir: /flashinfer configfile: pytest.ini collected 3240 items ... FAILED tests/gemm/test_mm_fp4.py::test_mm_fp4[mxfp4_alpha-False-True-cudnn-res_dtype1-512-512-256] - cudnn._compiled_module.cudnnGraphNotSupportedError: No valid engine configs for Matmul_MUL_ FAILED tests/gemm/test_mm_fp4.py::test_mm_fp4[mxfp4_alpha-False-True-cudnn-res_dtype1-512-512-512] - cudnn._compiled_module.cudnnGraphNotSupportedError: No valid engine configs for Matmul_MUL_ ================================================================================================================================ 180 failed, 270 passed, 2790 skipped, 1 warning in 8.97s ================================================================================================================================= ``` cudnn 9.14 --> unit test passes with 450 passed, 2790 skipped, 1 warning in 5.37s ``` =================================================================================================================================================== test session starts =================================================================================================================================================== platform linux -- Python 3.12.11, pytest-8.4.2, pluggy-1.6.0 rootdir: /flashinfer configfile: pytest.ini collected 3240 items tests/gemm/test_mm_fp4.py ... ====================================================================================================================================== 450 passed, 2790 skipped, 1 warning in 5.37s ======================================================================================================================================= ```   Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.). <!-- Optional: anything you'd like reviewers to focus on, concerns, e…

bkryu and others added 16 commits November 12, 2025 20:17

GLM v4 routing support

5c32462

formatting

fbeb950

fix rebase

05006af

.

a1df1d6

remove extra licenses

d170047

.

677fc59

.

814889a

murphymatt force-pushed the fix-kernel-launch-config branch from 73a9306 to 814889a Compare November 16, 2025 21:24

rebase on main

9d2226f

murphymatt requested a review from a team November 16, 2025 21:48

add tests

1e2dfbb

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Comments

Fix kernel launch config#5

Fix kernel launch config#5
murphymatt wants to merge 18 commits intomainfrom
fix-kernel-launch-config

murphymatt commented Nov 13, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

10 participants

Comments

Conversation

murphymatt commented Nov 13, 2025

📌 Description

🔍 Related Issues

🚀 Pull Request Checklist

✅ Pre-commit Checks

🧪 Tests

Reviewer Notes

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

10 participants