SM 107 Reland + Merge Back from v0.6.16 Release Branch - #4280
Merged
Merged
Conversation
This reverts commit 6258e52.
<!-- .github/pull_request_template.md --> ## 📌 Description The BMM & GEMM trtllm-gen cubins in ToT fail on SM 100. This MR introduces a WAR, loading an older set of published cubins when targeting SM 100. ## 🔍 Related Issues <!-- Link any related issues here --> ## 🚀 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 <!-- Optional: anything you'd like reviewers to focus on, concerns, etc. --> --------- Co-authored-by: Vinnie6167 <Vinnie6167@users.noreply.github.com> (cherry picked from commit 82b329b)
…lashinfer-ai#4189) ## 📌 Description Fixes the B300 sampling performance regression tracked in **NVBug 6517769**, targeting `release-v0.6.16` directly. [flashinfer-ai#4122](flashinfer-ai#4122) annotated 11 sampling kernels with `__launch_bounds__(BLOCK_THREADS)`. That annotation is effectively `__launch_bounds__(1024)` on every cc ≥ 8, which forces ptxas under a 64-reg/thread ceiling and causes measurable spills / slowdowns on B300 and H100. ### Why not a blanket removal? The annotations were not unexplained noise. Internal GitLab [MR !611](https://gitlab-master.nvidia.com/dl/flashinfer/flashinfer/-/merge_requests/611) (`Fix Rubin (SM 10.7) launch failures in sampling kernels via __launch_bounds__`, merged into `feat_sm107` 2026-04-28) documented the Rubin failure mode: > The Top-K sampling kernel was using >64 regs per thread with 1024 threads per block, exceeding the 65,536 register budget. The [Rubin Kernel Requirements Tracker](https://docs.google.com/spreadsheets/d/1Cn0_F6n_HMrzXPYreYbVIWbKQmfp-zzKne3Lz7akysU) recorded the pre-fix failures as `too many resources requested for launch` on `tests/utils/test_sampling.py` / logits sampling. That rationale never made the public flashinfer-ai#4122 description, but the release branch still carries SM107 support, so a full drop would re-break Rubin native launches. ### Fix Gate the annotation to native SM107 compiles only: ```cuda #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 1070) #define FLASHINFER_SAMPLING_LAUNCH_BOUNDS(block_threads) __launch_bounds__(block_threads) #else #define FLASHINFER_SAMPLING_LAUNCH_BOUNDS(block_threads) #endif ``` All 11 kernels use `FLASHINFER_SAMPLING_LAUNCH_BOUNDS(BLOCK_THREADS)` instead of unconditional `__launch_bounds__(BLOCK_THREADS)`. When today's JIT maps SM107 → `sm_100f` (`map_sm107_to_100f=True` while CUTLASS lacks `Sm107`), `__CUDA_ARCH__` is 1000 and the gate is inactive; that path inherits sm_100 register counts, which already fit a 1024-thread launch on the arches we measured. The gate is the safety net for **native `sm_107a`** compiles — the path MR !611 actually fixed. ## 🔍 Related Issues - NVBug 6517769 — B300 sampling perf drop after 290c091 - Internal MR !611 / commit `a8c10c24` on `feat_sm107` (Rubin launch-failure root cause) - Caused by flashinfer-ai#4122; flashinfer-ai#4171 reverted flashinfer-ai#4122 wholesale on `main` (no cherry-pick possible) ## 🧪 Tests ### ptxas (CUDA 13.0) with the gated macro | Target | Kernel | Result | |---|---|---| | sm_90a | MinP / TopK / TopP / TopKTopP | ≤48 regs, **no spills** (matches unconstrained) | | sm_103a | same | ≤48 regs, **no spills** | | sm_107a | — | **cannot compile** on this toolkit (`Unsupported gpu architecture`) | ### Wall-clock (NVBug repro on H100 NVL) ``` python3 flashinfer_benchmark.py --routine min_p_sampling_from_probs \ --batch_size 32 --vocab_size 129280 --min_p 0.1 -vv ``` | Config | median | |---|---| | release (unconditional `__launch_bounds__`) | 0.051 ms | | this PR (SM107-gated) | **0.045 ms** (same as full removal) | Correctness: earlier full-removal run of `tests/utils/test_sampling.py` on this header family was 1320 passed / 0 failed; the gated form is a no-op on H100 device code, so that result still applies. ## Reviewer Notes - Please confirm on **B300** that the NVBug case is restored, and on **Rubin** that `tests/utils/test_sampling.py` still launches when compiled for native `sm_107a` (and/or under the current `sm_100f` JIT mapping). - `flashinfer-ai#4171` already removed these annotations from `main` along with all SM107 enablement; this PR is release-only and preserves SM107 launchability. (cherry picked from commit 3436811)
…ild break (flashinfer-ai#4200) <!-- .github/pull_request_template.md --> Fixes both failing jobs in the 0.6.16rc2 release run ([30330417436](https://github.com/flashinfer-ai/flashinfer/actions/runs/30330417436)) on `c46207c2`. Both are deterministic consequences of introducing a second, Rubin-specific cubin pin (`TRTLLM_GEN_BMM_RUBIN` / `TRTLLM_GEN_GEMM_RUBIN`) alongside the existing one — neither is flaky. **1. `build-flashinfer-cubin` — `Failed to download cubins: checksum mismatch`** `get_checksums()` flattened every pin's `checksums.txt` into one dict keyed by **bare filename**, namespacing only `.h` files. The plain and Rubin pins ship the **same** sm100f/sm103a kernel names built from different sources, so the pin processed last silently overwrote the earlier one's hashes. `get_subdir_file_list()` then looked cubins up by bare name while headers already used the full path. Measured against the live published manifests: | pin pair | shared filenames | same hash | **conflicting** | |---|---|---|---| | `TRTLLM_GEN_BMM` vs `_RUBIN` | 2534 | 0 | **2534** | | `TRTLLM_GEN_GEMM` vs `_RUBIN` | 108 | 3 | **105** | The Rubin pins come later in `cubin_dirs`, so they win and 2640 of 8055 files are verified against the wrong pin's hash. Replaying both keying schemes over the real manifests: ``` OLD (bare filename) entries= 5412 files=8055 MISMATCHED=2640 NEW (full path) entries= 8055 files=8055 MISMATCHED=0 ``` **Fix:** key every entry by full path, and look cubins up the same way headers already were. This is **not** a stale pinned hash — all 12 `checksums.txt` manifests still match their `CheckSumHash` SHA256 exactly. It is also **not** download flakiness: the log shows two transient `403 Forbidden`s on `fmhaSm100fKernel_QkvE4m3OE2m1H256...Q{8,32}Kv128StaticSwapsAbForGen.cubin`, but both recovered on retry (no `Max retries reached`) and both URLs serve `200` now. A checksum-retry loop would have re-downloaded 2640 files, failed identically, and turned a fast failure into a slow one. **2. `build-flashinfer-jit-cache` (all 6 matrix entries) — compile error** ``` csrc/trtllm_batched_gemm_runner.cu(137): error: enum "batchedGemm::trtllm::gen::CudaArch" has no member "Sm107a" ``` The `Sm107a` comparison added in flashinfer-ai#4122 is unguarded, but that enumerator exists only in the Rubin pin's generated header — the default pin's `flashinferMetaInfo.h` has `Sm100f`/`Sm103a` and **zero** occurrences of `Sm107a` (the Rubin pin has 2552). So the non-Rubin module `fused_moe_trtllm_sm100` cannot compile. This is why `main`, where flashinfer-ai#4122 was reverted, is unaffected. **Fix:** guard with `TLLM_RUBIN_FEATURES`, which `gen_trtllm_gen_fused_moe_sm100_module` already defines (`rubin_flags`) exactly when the Rubin pin is selected — so the guard is by construction in sync with the enum's availability. `sm_version == 107` is unreachable in the non-Rubin module regardless: `enable_rubin` is chosen from the same device compute capability that `getSMVersion()` reports. **3. `tests/test_artifacts.py::test_get_subdir_file_list` — already broken on the release branch** Independent of the above, flashinfer-ai#4191 added the two Rubin pins to `cubin_dirs` without registering mocks for them, so the test exhausted its retry budget (~50 s of backoff) and died with `FileNotFoundError`. Verified failing at the unmodified base `c46207c2`. Mocks added, meta-info header count corrected 3 → 5, plus regression coverage for the collision above. - Release run: https://github.com/flashinfer-ai/flashinfer/actions/runs/30330417436 - Introduced by the per-arch pin split (flashinfer-ai#4191) interacting with SM107 support (flashinfer-ai#4122) 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/test_artifacts.py`: **1 failed, 3 passed → 4 passed**. The new regression assertion is verified to actually catch the bug — reverting only `artifacts.py` to the buggy base while keeping the new test yields: ``` AssertionError: Bmm_..._schedS_bN_clmp_dynBatch_sm100f.cubin resolved to the same checksum for both pins (bbbb111122223333) -- the per-pin hashes collided ``` - **Scope of local testing.** `tests/test_artifacts.py` is fully exercised locally (it is `responses`-mocked, no GPU needed). The two build jobs themselves are not reproducible locally and need CI to confirm. - **Why the jit-cache fix is sufficient for the whole matrix.** All 4 failed jit-cache jobs report **exactly one** compiler error each — the same `Sm107a` one, same target `csrc_trtllm_batched_gemm_runner.cuda.o`. The matrix builds `7.5 8.0 8.9 9.0a 10.0a 10.3a 12.0f` (no 10.7), so `has_sm107` is false and the Rubin modules are not compiled there — this change unblocks the job without pulling in any `sm_107a` codegen. - **Two alternatives considered and rejected** for the `Sm107a` guard: - A new `FLASHINFER_HAS_SM107A` macro — nothing would define it, so the `#else` branch would always compile, silently disabling Rubin kernel selection. - Falling back to `Sm103a`/`Sm100f` under `sm_version == 107` — worse than the build error: on real Rubin silicon it would select wrong-arch cubins instead of failing loudly. - **Follow-up worth considering (not in this PR):** the `checksum mismatch` error does not name the offending file. Including the path plus expected/actual hashes would have made this diagnosable from the log alone. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude Opus 5 <noreply@anthropic.com> (cherry picked from commit 6327bb8)
…ai#4215) ## 📌 Description `gen_moe_utils_module()` is the only MoE JitSpec that does not pass `map_sm107_to_100f=True`. On Rubin (SM107) this breaks the build outright: ``` nvcc fatal : Unsupported gpu architecture 'compute_107a' ``` No public CUDA toolkit (13.0 or 13.2) accepts `compute_107a`, so the SM107 → `sm100f` family mapping is required. ### Why the omission is not merely redundant These flags are passed as `extra_cuda_cflags`. In `build_cuda_cflags` (`flashinfer/jit/cpp_ext.py`), the `module_has_gencode` branch lets module-level flags **override** the globally computed, correctly-mapped flags. So leaving the argument off does not fall back to the global mapping — it *removes* it. ### Blast radius Strictly a no-op off Rubin. `CompilationContext.get_nvcc_flags_list` only rewrites the gencode when `(major, minor) == (10, "7a")`: ```python if apply_sm107_mapping and major == 10 and minor == "7a": flags.append("-gencode=arch=compute_100f,code=sm_100f") ``` `apply_sm107_mapping` is additionally gated on `not cutlass_supports_sm107()`, so once the bundled CUTLASS gains native `compute_107a` support this stops applying automatically, with no further code change. Affects the CuteDSL MoE path — `gen_moe_utils_module()` is reached from `aot.py` and `flashinfer/fused_moe/cute_dsl/moe_utils.py`. ## 🔍 Related Issues Found while investigating the Rubin NVFP4 MoE failure. This is **independent of** and complementary to flashinfer-ai#4168 / flashinfer-ai#4213 (incompatible output-scale cubins) — that one fixes a runtime correctness bug, this one fixes a build break. Both are needed on Rubin. ## 🚀 Pull Request Checklist ### ✅ 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. ## 🧪 Tests - [x] All tests are passing (`unittest`, etc.). Verified on Hecate GR100 (cc 10.7, aarch64): without this change the `moe_utils` JIT module fails to compile on SM107; with it the module builds and the MoE suites run. No test added — this is a build-configuration fix on a JitSpec, exercised by any SM107 build of the CuteDSL MoE path. ## Reviewer Notes Single-line change plus comment. Every other MoE JitSpec already opts into this mapping; this one appears to have simply been missed when SM107 support was added in flashinfer-ai#4122. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com> (cherry picked from commit d4635e3)
…k race test (flashinfer-ai#4225) ## Summary Fixes two test regressions introduced by flashinfer-ai#4191 (release-only, not on `main`): - **flashinfer-ai#4208** — `test_make_tuning_config_reuses_topk_ids_initializer` called `get_trtllm_moe_sm100_module.cache_clear()`, but flashinfer-ai#4191 moved `@functools.cache` to `_get_trtllm_moe_sm100_module_impl`. - **flashinfer-ai#4166 (re-open)** — flashinfer-ai#4187 updated the symlink test to assert under `GEN_SRC_DIR/flashinfer/...`, but flashinfer-ai#4191 moved the export symlink to `GEN_SRC_DIR/trtllm_export/fused_moe_trtllm_sm100/flashinfer/...`. Both failures appeared on all Blackwell jobs in Test Cycle 2 (pipeline 59912149, rc3). ## Test plan - [x] `pytest tests/autotuner/test_autotuner_core.py::test_make_tuning_config_reuses_topk_ids_initializer` — pass - [x] `pytest tests/utils/test_gen_module_symlink_race_condition.py` — pass - [ ] B300/GB200 CI rerun to confirm cycle-2 green on these two files (cherry picked from commit 95c3ec1)
…(release port for flashinfer-ai#4107) (flashinfer-ai#4230) Release-specific port of flashinfer-ai#4177 onto `release-v0.6.16` for flashinfer-ai#4107. On SM12x (Spark, RTX Pro 6000), TRTLLM routed-MoE backends were incorrectly claiming support (`arch >= 100`) and then dispatching sm100f/sm103a cubins, causing `RuntimeError: Error occurred when running GEMM!` or segfaults in `test_split_fused_moe_kernel_vs_reference`. - **`csrc/trtllm_batched_gemm_runner.cu`**: Replace per-SM if-chains with `isArchCompatible()`; reject unknown cubin families; guard `Sm107a` behind `#ifdef TLLM_RUBIN_FEATURES` (only exists in the Rubin cubin pin's headers). - **`csrc/trtllm_gemm_runner.cu`**: Same arch filter for the plain GEMM runner (previously had no arch filtering at all). - **`flashinfer/fused_moe/api.py`**: Tighten `Trtllm*Config.supported()` from `arch >= 100` to explicit allowlists `_TRTLLM_ROUTED_ARCHS = (100, 103, 107)` and `_TRTLLM_ROUTED_FP8_ARCHS = (100, 103)`. - **`tests/moe_ep/test_split_fused_moe_kernel_vs_reference.py`**: Gate GPU tests on `config_cls.supported(arch)`; add CPU contract tests + SM120 regression guard. - Includes **sm107** in the allowlist (unlike the rebased flashinfer-ai#4177 for `main`, which will drop 107 after the flashinfer-ai#4122 revert). - `Sm107a` enum case is gated behind `TLLM_RUBIN_FEATURES`, matching the existing release pattern from flashinfer-ai#4122/flashinfer-ai#4191. Verified against the actual pinned headers: the default BMM/GEMM pins do not define `Sm107a`; only the Rubin pins do. - CPU tests: 36 passed, 2 skipped (`test_split_fused_moe_kernel_vs_reference.py`) - Backend claims: sm120/sm121 now fall back to Cutlass only (no Trtllm* backends) - Compile: `isArchCompatible()` builds cleanly against both default and Rubin BMM export headers; ported batched runner compiles against default pin - Fixes flashinfer-ai#4107 - Upstream PR: flashinfer-ai#4177 (targets `main`, has merge conflicts) - Cherry-pick source commit on `main`: not yet merged `flashinfer-ai#4213` on release references `options.mDtypeSfC`, which does not exist in the default BMM cubin pin's headers (only in the Rubin pin). This is a separate release-only compile issue on the non-Rubin module, predating this port. (cherry picked from commit 29d360f)
…ackages carry mDtypeSfC (flashinfer-ai#4235) ## 📌 Description Fixes two coupled problems on `release-v0.6.16`: **1. The branch tip cannot compile the trtllm-gen fused-MoE module — on any architecture.** flashinfer-ai#4213 checks `BatchedGemmOptions::mDtypeSfC`, but that field does not exist in either pinned cubin package's exported headers (Blackwell `b368d003`, Rubin `46d3f356` — verified 0 occurrences in both published `BatchedGemmOptions.h`, and in the internal `81a53cff` package as well). Any JIT or AOT rebuild of `fused_moe_trtllm_*` fails with: ``` trtllm_batched_gemm_runner.cu(128): error: class "BatchedGemmOptions" has no member "mDtypeSfC" ``` (The original PR flashinfer-ai#4168 shows the same failure in its CI on main.) **2. The FP4-MoE illegal-memory-access flashinfer-ai#4213 targets (flashinfer-ai#4164 / NVBug 6517914) is still live on Rubin.** The offending `bmm_E2m1xFp32_*` kernels emit linear FP32 output scale-factors into the E4M3-sized (1 B/block) buffer the MoE pipeline allocates → device-side overflow → async IMA, near-deterministic under autotune. Verified against the published catalogs: **132 such kernels, all in the sm_107a package (`46d3f356`); zero in the Blackwell package** — which is why the per-arch pin split (flashinfer-ai#4191) healed Blackwell while Rubin still crashes. ## What changed Replace the un-compilable typed check with a kernel-**name** filter carrying the same intent: reject configs whose `mFunctionName` contains the `E2m1xFp32` token. This compiles against the *currently published* packages and removes the bad candidates before autotune can run them. flashinfer-ai#4213's other changes (per-token-scaling restructure, launcher tactic-enumeration alignment) are inherited unchanged. **Scope note:** the name filter is deliberately narrower than the typed check — it rejects only the known-bad family. Against the published catalogs this is behaviorally equivalent (the `E2m1xFp32` kernels are the only offenders present). The TODO in the code ties restoring the typed `mDtypeSfC` check to publishing + pinning packages that carry the field. ## Validation - Build break reproduced (sbsa AOT wheel build, CTK 13.4) at the branch tip; gone with this patch. - VR200 (Rubin) re-run of the affected MoE test files in progress; results will be posted here before marking ready for review. Refs: flashinfer-ai#4164, flashinfer-ai#4213, flashinfer-ai#4168 · NVBug 6517914 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-authored-by: Vinnie6167 <Vinnie6167@users.noreply.github.com> (cherry picked from commit b9bb507)
…flashinfer-ai#4226) <!-- .github/pull_request_template.md --> ## 📌 Description <!-- What does this PR do? Briefly describe the changes and why they’re needed. --> The CuTe DSL support check did not successfully set the compilation target to SM 100f when the DSL did not support SM107 on Rubin. Setting the env variable CUTE_DSL_ARCH after initializing CuTe DSL does not change the arch target. Thus, this PR changes the gating to checking the CUTE_DSL_ARCH env is set properly, instead of setting it and returning true. This PR also adds the CuTe DSL support checks to the norm CuTe DSL kernels. ## 🔍 Related Issues <!-- Link any related issues here --> ## 🚀 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 <!-- Optional: anything you'd like reviewers to focus on, concerns, etc. --> --------- Co-authored-by: Vinnie6167 <Vinnie6167@users.noreply.github.com> (cherry picked from commit d78020f)
## What
Adds a `skip_if_nvfp4_kv_unsupported()` helper to
`tests/attention/test_batch_decode_kernels.py` and calls it at the top
of `test_batch_decode_with_paged_kv_cache_nvfp4`, skipping the NVFP4-KV
decode tests on SM107 (compute capability `(10, 7)`).
## Why
No NVFP4-KV (E2m1) decode kernels exist for SM107: the published
trtllm-gen FMHA package contains E2m1-KV kernels only for Sm100a and
Sm103a (264 each, fp8-out only) — zero Sm100f or Sm107a — and sm_100a
binaries cannot load on sm_107.
The library already guards this deliberately (`flashinfer/decode.py`,
`flashinfer/prefill.py` raise `ValueError("KV Cache NVFP4 is not
supported on SM107")`), so on Rubin hardware all 289 nvfp4 ids in this
file currently **fail** with that exact ValueError instead of skipping:
```
[289 failed, 2970 passed] ValueError: KV Cache NVFP4 is not supported on SM107 (x289)
```
(full release-v0.6.16 sweep on VR200, 2026-07-28)
## Notes
- The skip keys on `== (10, 7)` to mirror the library guard exactly — it
skips precisely where the library declares unsupported, nowhere else.
- All nvfp4 entry points route through
`test_batch_decode_with_paged_kv_cache_nvfp4` (including the
`_large_head` variant and the torch-compile smoke test), so the single
call site covers all 289 ids.
- When NVFP4-KV kernels are generated for SM107/SM100f, lifting this
skip belongs to the same change that lifts the library guard.
🤖 Generated with [Claude Code](https://claude.com/claude-code)
Co-authored-by: Vinnie6167 <Vinnie6167@users.noreply.github.com>
(cherry picked from commit efd289a)
…ai#4258) ## Summary Adds a `DEEPGEMM_RUBIN` entry to the artifact metadata in `flashinfer/artifacts.py`, mirroring the existing non-RUBIN `DEEPGEMM` var: - **`ArtifactPath.DEEPGEMM_RUBIN`** → `7ec7ac40b9fd48172651b77ff2ebe20d79decc39/deep-gemm/` - **`CheckSumHash.DEEPGEMM_RUBIN`** → `09e961d4e3852a6cf81b3482d0604c09dcb1f69c1b7936f535c9ee2f53335184` - Wired the checksum into `CheckSumHash.map_checksums` via `safe_urljoin(ArtifactPath.DEEPGEMM_RUBIN, "checksums.txt")`. Targets the `release-v0.6.16` release branch. ## Notes Scope is limited to the path + checksum as requested. Unlike the `TRTLLM_GEN_*_RUBIN` pins, `DEEPGEMM_RUBIN` is **not** added to `get_subdir_file_list()`'s `cubin_dirs`, so it is not fetched by `flashinfer artifacts download` — let me know if that wiring is also wanted. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com> (cherry picked from commit f07ea8f)
…4261) ## Summary Follow-up to flashinfer-ai#4258 (which added `ArtifactPath.DEEPGEMM_RUBIN` and its checksum). That PR merged the path + checksum only; this PR adds the arch-based split that actually routes deep-gemm cubin loading to the Rubin directory on sm_107, mirroring the `enable_rubin` pattern from flashinfer-ai#4191. ### `flashinfer/deep_gemm.py` - New `is_rubin_arch()` / `get_deepgemm_artifact_path()` helpers selecting `DEEPGEMM_RUBIN` on arch `107a` and `DEEPGEMM` otherwise. - Used in `load_all()`, `load()`, and `KernelMap.init_indices()` so both the kernel cubins and `kernel_map.json` are fetched from the correct per-arch directory. - Added `KernelMap.KERNEL_MAP_HASH_RUBIN = f8bf2b1b…a36e27` for the Rubin `kernel_map.json` (a separate manifest from the default one), selected by arch. ### `flashinfer/artifacts.py` - Added `DEEPGEMM_RUBIN` to `get_subdir_file_list()`'s download list so `flashinfer artifacts download` fetches it, and updated the DEEPGEMM comment to reference the Rubin map hash. Both Rubin hashes were verified by downloading the artifacts directly from the cubin repository. Targets the `release-v0.6.16` release branch. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com> (cherry picked from commit d4c14f1)
Rubin-specific part of flashinfer-ai#4252's conflict resolution (c4053a2; the flashinfer-ai#4180 base it carried is already on main as a34a735). The pinned Rubin BMM package (TRTLLM_GEN_BMM_RUBIN) predates SiTuGlu and has no 192-tile FP4 kernels, so both are compiled out under -DTLLM_RUBIN_FEATURES (set only for the sm_107 module). (cherry picked from commit c4053a2828e176d1a99f172a2a71a8fdcbcbe10c, Rubin guards only)
These asserts were written on main after the flashinfer-ai#4171 revert (via flashinfer-ai#4159) and lock in 'no SM107'. With SM107 support re-landed, TrtllmFp4Config and TrtllmBf16Config claim 107 again (_TRTLLM_ROUTED_ARCHS); the FP8 backends correctly remain SM100/SM103-only.
Contributor
|
Caution The consumer version of Gemini Code Assist on GitHub has been sunset. All code review activity has officially ceased. |
Contributor
📝 WalkthroughWalkthroughSM107/Rubin support is added across compilation, AOT/JIT module generation, artifact loading, TensorRT-LLM kernel dispatch, backend capability checks, CuTe DSL validation, NVFP4 restrictions, and architecture-aware tests. ChangesSM107 and Rubin build pipeline
Estimated code review effort: 5 (Critical) | ~120 minutes Possibly related PRs
Suggested labels: Suggested reviewers: 🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✨ Finishing Touches 💡 1🛠️ Fix failing CI checks 💡
🧪 Generate unit tests (beta)
Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out. Comment |
Vinnie6167
marked this pull request as ready for review
July 30, 2026 22:11
Contributor
|
Caution The consumer version of Gemini Code Assist on GitHub has been sunset. All code review activity has officially ceased. |
Vinnie6167
requested review from
bkryu,
ishovkun,
jiahanc,
jimmyzho,
kahyunnam,
yongwww and
yzh119
as code owners
July 30, 2026 22:11
This was referenced Jul 31, 2026
5 tasks
5 tasks
bkryu
added a commit
that referenced
this pull request
Aug 1, 2026
…s + CuTe-DSL MoE device guard) (#4301) <!-- .github/pull_request_template.md --> ## 📌 Description <!-- What does this PR do? Briefly describe the changes and why they’re needed. --> `main` is currently red on multiple test files, all stemming from #4280. This PR bundles the fixes so CI can go green in a single run — the gate requires one passing run, and none of these fixes pass in isolation (each file still fails on the others), so they cannot land separately. **Supersedes #4292** . Current PR includes the commit from #4292 and adds the remaining MoE fix on top to pass the CI ## 🔍 Related Issues <!-- Link any related issues here --> - #4280 - #4292 ## 🚀 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 <!-- Optional: anything you'd like reviewers to focus on, concerns, etc. --> <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **Bug Fixes** * Improved error messages when checksum manifests cannot be downloaded and are unavailable locally, including details about the affected artifact. * Preserved the use of cached checksum manifests when refresh attempts fail, preventing unnecessary errors when valid local data is available. * Improved reliability when validating artifacts across multiple hardware and software configurations. <!-- end of auto-generated comment: release notes by coderabbit.ai --> --------- Co-authored-by: Alex Yang <aleyang@nvidia.com> Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
aleozlx
added a commit
that referenced
this pull request
Aug 4, 2026
…s + CuTe-DSL MoE device guard) (#4301) <!-- .github/pull_request_template.md --> ## 📌 Description <!-- What does this PR do? Briefly describe the changes and why they’re needed. --> `main` is currently red on multiple test files, all stemming from #4280. This PR bundles the fixes so CI can go green in a single run — the gate requires one passing run, and none of these fixes pass in isolation (each file still fails on the others), so they cannot land separately. **Supersedes #4292** . Current PR includes the commit from #4292 and adds the remaining MoE fix on top to pass the CI ## 🔍 Related Issues <!-- Link any related issues here --> - #4280 - #4292 ## 🚀 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 <!-- Optional: anything you'd like reviewers to focus on, concerns, etc. --> <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **Bug Fixes** * Improved error messages when checksum manifests cannot be downloaded and are unavailable locally, including details about the affected artifact. * Preserved the use of cached checksum manifests when refresh attempts fail, preventing unnecessary errors when valid local data is available. * Improved reliability when validating artifacts across multiple hardware and software configurations. <!-- end of auto-generated comment: release notes by coderabbit.ai --> --------- Co-authored-by: Alex Yang <aleyang@nvidia.com> Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com> (cherry picked from commit d020372)
aleozlx
added a commit
that referenced
this pull request
Aug 4, 2026
Mirrors the v0.6.16 revert (#4246, commit 0a27ba3) so that 0.6.17 does not ship the API-breaking ring-buffer cache contract that 0.6.16 left out. Reverts: afd4754 mamba checkpointing SSU: two-kernel split + ring-buffer cache for checkpointing SSU (#3975) f90e9c4 docs(mamba): document checkpointing varlen arguments (#4129) Unlike release-v0.6.16, this branch also carries #4129, so both are reverted here; 0.6.16 only needed #3975. Verified no collateral damage: the SM107 changes from #4280 in tests/mamba/conftest.py and the #4029 changes in flashinfer/utils.py are preserved. NOTE: like the 0.6.16 revert, this also reverts the is_cvt_rs_supported correctness fix that rode along in #3975 (back to `major in (10, 11)`). See the release notes discussion -- that hunk is a candidate to keep. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
This was referenced Aug 5, 2026
5 tasks
aleozlx
added a commit
that referenced
this pull request
Aug 7, 2026
#4080) on release-v0.6.17 (#4411) ## 📌 Description Reverts #3738 (`5823159c`) and its two dependents on `release-v0.6.17`, to unblock the release after a **critical crash reported by a vLLM user that escaped QA**. This is a release-unblocking revert, not a judgement on the feature. The proper fix is expected in **0.6.18**. The same revert is being applied to `release-v0.6.16` separately by another engineer — note #3738 shipped in v0.6.16 and v0.6.16.post2, so 0.6.17 is not the first release carrying it. ## Why three commits #3738 cannot be reverted alone — two later commits are built on it: | commit | PR | why it must go first | |---|---|---| | `2475121f` | #4080 | fixes `getProfilerWorkspaces` in code #3738 reworked (the `quant_1 && … && quant_6` assert) | | `c83607a9` | #4025 | exists *only* to re-export the interleave helpers that #3738 moved `core` → `prepare` | | `5823159c` | #3738 | the target | Reverted newest-first. #4080 and #4025 reverted with **zero conflicts**; #3738 conflicted in 3 files. ## Conflict resolution — later work is preserved This reverts #3738 only, **not** the features layered on top of it: - **`core.py`** — kept `profile_ids` / `workspace_buffer` and the CUDA-device guard (#4057 and follow-ups); dropped only `use_wfp4afp8_humming` and its plumbing. - **`__init__.py`** — kept `cutlass_fused_moe_workspace_size` (#4057); restored the two `interleave_moe_*_for_sm90_mixed_gemm` exports, which move back to `core` once #3738 is gone. - **`prepare.py`** — dropped the SM90 Humming docstring and the then-unused `functools` / `struct` imports; kept the `TrtllmBf16Config` wording added after #3738. Net: **47 files, +1701 / -10557** — a near-exact mirror of #3738's +10546 / -1704, the delta being the later features deliberately kept. ## ✅ Verification - No `wfp4afp8_humming` / `Humming` reference remains anywhere in `flashinfer/`, `csrc/` or `tests/`. - **Every** `cutlass_kernels/` and `cutlass_backend/` source matches the pre-#3738 parent **except two files**, and both residuals are fully attributable to later commits that correctly survive: - `flashinfer_cutlass_fused_moe_binding.cu` → #4057 (caller-owned workspace; 16 workspace-API references retained) - one `std::remove_reference_t<…>` line in `moe_gemm_tma_ws_launcher.inl` → SM107 work (#4122 / #4280) - Every relative import in `fused_moe/__init__.py` resolves (AST-checked). - `py_compile` passes on all touched Python files. - `pre-commit` (clang-format, mypy, ruff check, ruff format) passes. **✅ GPU-validated** — see [the validation comment](#4411 (comment)). A/B contrast on **B300 / sm103** with the reporter's repro, same node / container / install, only the commit differing: | | commit | result | |---|---|---| | baseline | `4e1206b3` (branch head, unreverted) | `FAIL_IN_AUTOTUNE` — dies at gemm1 tactic 1/21 | | revert | `57d74695` (this PR) | **`PASS`** — gemm1 21/21, gemm2 21/21 | The `without autotune` control passes on both builds, isolating the autotuned path. **Still not run:** the MoE unit-test suites. CI remains the gate for regression coverage; this validates the specific reported crash only. ## Reviewer notes - Worth a targeted look at the `core.py` conflict resolutions, since that is where #3738 and #4057 interleaved in the same parameter lists. - If the crash turns out to be reproducible on v0.6.16 too, that confirms #3738 as the cause rather than a 0.6.17-specific interaction. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
This was referenced Aug 10, 2026
jimmyzho
added a commit
that referenced
this pull request
Aug 18, 2026
) #4375 CU_FUNC_ATTRIBUTE_SHARED_MEMORY_MODE, CU_SHARED_MEMORY_MODE_ALLOW_OVERSIZED_SHARED_MEMORY and CU_DEVICE_ATTRIBUTE_MAX_OVERSIZED_SHARED_MEMORY_PER_BLOCK were introduced in CUDA 13.4; CUDA 13.3 only ships the launch-attribute half of the shared-memory-mode API (CU_LAUNCH_ATTRIBUTE_SHARED_MEMORY_MODE, CUsharedMemoryMode values 0-2). The `#if CUDA_VERSION >= 13030` guards added in #4122 / relanded in #4280 therefore break the trtllm-gen fmha_gen JIT build on CUDA 13.3 (the current released toolkit) with fmhaKernels.cuh(136/157/169): error: identifier ... is undefined on any GPU arch, e.g. via trtllm_batch_decode_with_kv_cache_mla on SM100. Bump the three guards to `#if CUDA_VERSION >= 13040`. No CUDA 13.3 fallback is needed: the 13.3 driver does not support the oversized mode, so falling through to the existing MAX_DYNAMIC_SHARED_SIZE_BYTES path is correct. References: - CUDA 13.3 Driver API (symbols absent): https://docs.nvidia.com/cuda/archive/13.3.0/cuda-driver-api/group__CUDA__TYPES.html - CUDA 13.4 developer preview Driver API (symbols present): https://docs.nvidia.com/cuda/developer-preview/13.4/pdf/CUDA_Driver_API.pdf Verified with CUDA 13.3 (nvcc V13.3.73) on SM100: tests/attention/test_trtllm_gen_mla.py::test_trtllm_batch_decode_mla (trtllm-gen backend) fails to build before, passes after. <!-- .github/pull_request_template.md --> ## 📌 Description <!-- What does this PR do? Briefly describe the changes and why they’re needed. --> ## 🔍 Related Issues <!-- Link any related issues here --> ## 🚀 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 <!-- Optional: anything you'd like reviewers to focus on, concerns, etc. --> <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **Bug Fixes** * Updated CUDA compatibility checks for oversized shared-memory support. * Ensured related functionality is available only with CUDA 13.4 or newer, improving compatibility with supported CUDA environments. <!-- end of auto-generated comment: release notes by coderabbit.ai --> Co-authored-by: Claude Fable 5 <noreply@anthropic.com> Co-authored-by: Jimmy Zhou <79552142+jimmyzho@users.noreply.github.com>
jefby
pushed a commit
to jefby/flashinfer
that referenced
this pull request
Aug 19, 2026
…ashinfer-ai#4377) flashinfer-ai#4375 CU_FUNC_ATTRIBUTE_SHARED_MEMORY_MODE, CU_SHARED_MEMORY_MODE_ALLOW_OVERSIZED_SHARED_MEMORY and CU_DEVICE_ATTRIBUTE_MAX_OVERSIZED_SHARED_MEMORY_PER_BLOCK were introduced in CUDA 13.4; CUDA 13.3 only ships the launch-attribute half of the shared-memory-mode API (CU_LAUNCH_ATTRIBUTE_SHARED_MEMORY_MODE, CUsharedMemoryMode values 0-2). The `#if CUDA_VERSION >= 13030` guards added in flashinfer-ai#4122 / relanded in flashinfer-ai#4280 therefore break the trtllm-gen fmha_gen JIT build on CUDA 13.3 (the current released toolkit) with fmhaKernels.cuh(136/157/169): error: identifier ... is undefined on any GPU arch, e.g. via trtllm_batch_decode_with_kv_cache_mla on SM100. Bump the three guards to `#if CUDA_VERSION >= 13040`. No CUDA 13.3 fallback is needed: the 13.3 driver does not support the oversized mode, so falling through to the existing MAX_DYNAMIC_SHARED_SIZE_BYTES path is correct. References: - CUDA 13.3 Driver API (symbols absent): https://docs.nvidia.com/cuda/archive/13.3.0/cuda-driver-api/group__CUDA__TYPES.html - CUDA 13.4 developer preview Driver API (symbols present): https://docs.nvidia.com/cuda/developer-preview/13.4/pdf/CUDA_Driver_API.pdf Verified with CUDA 13.3 (nvcc V13.3.73) on SM100: tests/attention/test_trtllm_gen_mla.py::test_trtllm_batch_decode_mla (trtllm-gen backend) fails to build before, passes after. <!-- .github/pull_request_template.md --> ## 📌 Description <!-- What does this PR do? Briefly describe the changes and why they’re needed. --> ## 🔍 Related Issues <!-- Link any related issues here --> ## 🚀 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 <!-- Optional: anything you'd like reviewers to focus on, concerns, etc. --> <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **Bug Fixes** * Updated CUDA compatibility checks for oversized shared-memory support. * Ensured related functionality is available only with CUDA 13.4 or newer, improving compatibility with supported CUDA environments. <!-- end of auto-generated comment: release notes by coderabbit.ai --> Co-authored-by: Claude Fable 5 <noreply@anthropic.com> Co-authored-by: Jimmy Zhou <79552142+jimmyzho@users.noreply.github.com>
kahyunnam
added a commit
that referenced
this pull request
Aug 19, 2026
Mirrors the v0.6.16 revert (#4246, commit 0a27ba3) and its v0.6.17 counterpart (09fd5fc) so that 0.6.18 does not ship the API-breaking ring-buffer cache contract that neither 0.6.16 nor 0.6.17 shipped. Reverts: afd4754 mamba checkpointing SSU: two-kernel split + ring-buffer cache for checkpointing SSU (#3975) f90e9c4 docs(mamba): document checkpointing varlen arguments (#4129) Like release-v0.6.17, this branch carries #4129 as well, so both are reverted; 0.6.16 only needed #3975. main still carries both. Verified no collateral damage: the SM107 change from #4280 in tests/mamba/conftest.py and the #4029/#4078 changes in flashinfer/utils.py are preserved, and the five core reverted files now byte-match release-v0.6.17. NOTE: like both earlier reverts, this also reverts the cvt_rs fix that rode along in #3975 -- is_cvt_rs_supported goes back to `major in (10, 11)` (wrong for SM110a) and the CUDA guard back to SM100_ALL only (B300/sm_103a falls to software emulation). That matches what 0.6.17 shipped, but the hunk remains a candidate to keep.
Vinnie6167
added a commit
to Vinnie6167/flashinfer
that referenced
this pull request
Aug 21, 2026
…t it is_cute_dsl_arch_supported's docstring said the probe "pins the DSL's default target via the CUTE_DSL_ARCH environment variable". That was true when flashinfer-ai#4122 landed -- it called os.environ.setdefault("CUTE_DSL_ARCH", family) -- but the line went away with the flashinfer-ai#4171 revert and was not restored by the flashinfer-ai#4280 re-land, leaving the wording stale. The passive behaviour is the correct one, and FlashInfer should not set the variable: as the comment a few lines below notes, an env var set after cutlass is imported does not retarget the DSL, and the setdefault sat after `from cutlass.base_dsl.arch import Arch` so it was a no-op anyway. Retargeting is the user's call, since it affects every DSL kernel in the process. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
5 tasks
kahyunnam
added a commit
that referenced
this pull request
Aug 27, 2026
## Summary Fixes #4773: `mm_fp8` / `test_mm_fp8_replay` SIGSEGV in `cuModuleGetFunction` during low-latency GEMM autotune on SM100 (B200/GB200) and SM103 (B300). After #4648 the trtllm-gen GEMM pack is a single multi-arch publish, so `flashinferMetaInfo.h` now lists sm100f **and** sm107a configs in the manifest the Blackwell module compiles against. Previously the Rubin cubins lived in a separate `TRTLLM_GEN_GEMM_RUBIN` pin, so the non-Rubin manifest was sm100-only. `trtllm_low_latency_gemm_runner.cu` was the one trtllm-gen runner without the arch filter that #4280 added to its siblings. With the consolidated pack, `getValidTactics()` returned 16 tactics on Blackwell (8 sm100f + 8 sm107a); `cuModuleLoadData` fails on the first sm107a cubin, the `CUresult` is ignored by the generated `GemmInterface`, and `cuModuleGetFunction` faults on the uninitialised `CUmodule`. Two commits, no cubin regeneration: 1. **`isArchCompatible` filter** when building `mPassingConfigIndices`, identical to `csrc/trtllm_gemm_runner.cu` (`Sm107a` only under `TLLM_RUBIN_FEATURES`, `Sm100f` allowed on sm100 and sm103). 2. **`checkPassingConfigIndex` in `run()`**, also mirroring `trtllm_gemm_runner.cu`. Tactic ids are manifest indices, and the autotuner's file-config key (`custom_op`, `runner_class_name`, `nearest_profile`, `extras`) does not include the device arch, so a config saved via `save_configs()` / `autotune(cache=...)` on other hardware — or an explicit FFI tactic — could still hand a foreign-arch index straight to the cubin loader. It now raises instead of faulting. Other ops touched by #4648 already have the equivalent guard, so no further coverage is needed: | Consumer | Arch filter | |---|---| | `trtllm_gemm_runner.cu` | `isArchCompatible` + `checkPassingConfigIndex` (#4280) | | `trtllm_batched_gemm_runner.cu` (trtllm-gen MoE backend) | `isArchCompatible` + `checkPassingConfigIndex` (#4280) | | trtllm-gen FMHA | `isSMCompatible()` in `fmhaKernels.cuh`, with explicit sm107 rules | | `trtllm_low_latency_gemm_runner.cu` | **missing — this PR** | ## Test plan Local B200 (SM100, CUDA 13.0, Python 3.10), on `release-v0.6.18` + these commits: - [x] Before the filter: 16 valid tactics (sm100f indices `0,2,3,4,5,7,10,11` + sm107a `93,95,96,97,101,102,104,109`); SIGSEGV on the first sm107a cubin load. - [x] After: 8 sm100f tactics only; `mm_fp8` passes under `autotune()` and on the heuristic `tactic=-1` path. - [x] Forced sm107a tactic (`93`) now raises `RuntimeError: Tactic 93 is not in this runner's compatible config set` instead of SIGSEGV. - [x] `pytest tests/gemm/test_mm_fp8.py tests/utils/test_logging_replay.py` → 44 passed, 2 skipped (includes `test_mm_fp8_replay`, the test that crashed in CI). - [ ] GitLab `unit_test_b300` / GB200 jobs covering `tests/gemm/test_mm_fp8.py` and `tests/utils/test_logging_replay.py`.
5 tasks
kahyunnam
added a commit
that referenced
this pull request
Sep 1, 2026
## 📌 Description Port of #4786 (plus the #4792 `Sm100f`/sm107 allowance) onto `main`. After #4648 the trtllm-gen GEMM pack is a single multi-arch artifact, so `getValidTactics()` on the low-latency runner returned SM107 cubins on Blackwell. Autotune then handed those indices to `cuModuleLoadData`. `trtllm_low_latency_gemm_runner.cu` was the one trtllm-gen runner still missing the `isArchCompatible` / `checkPassingConfigIndex` filter that #4280 added to `trtllm_gemm_runner.cu` and `trtllm_batched_gemm_runner.cu`. Do not cherry-pick #4786 verbatim: that filter treated `Sm100f` as sm100/sm103 only, and `select_kernel()` still names `_sm100f` heuristics, which #4792 showed fails every `mm_fp8` case on Rubin. This PR uses the combined `release-v0.6.18` rule (`Sm100f` on 100/103/107). On B200 (SM100) unpatched `getValidTactics(4, 2560, 8192)` returned 16 indices (`0,2,3,4,5,7,10,11` + `93,95,96,97,101,102,104,109`). Forced tactic `93` failed inside `gemm.run`. After the filter: 8 sm100f tactics; tactic `93` raises `Tactic 93 is not in this runner's compatible config set`. ## 🔍 Related Issues - Closes #4773 - Cherry-pick / port of #4786 onto `main` (not a verbatim cherry-pick; includes the #4792 `Sm100f` sm107 allowance) ## 🚀 Pull Request Checklist ### ✅ Pre-commit Checks - [x] I have installed `pre-commit`. - [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. ## 🧪 Tests - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.). Local B200 (SM100, CUDA 13.0, torch 2.13.0+cu130), worktree at `upstream/main` + this commit: - `pytest tests/gemm/test_mm_fp8.py tests/utils/test_logging_replay.py::test_mm_fp8_replay` → 31 passed - Tactic dump and forced-sm107a FFI path as above Not verified here: SM103 (B300) or SM107 (Rubin). Those were covered on `release-v0.6.18` by #4786 / #4792. ## Reviewer Notes Sibling runners on `main` still map `Sm100f` to `100 || 103` only. They were left alone: they already have an arch filter, their heuristics have dedicated `_sm107a` names, and #4792 called that follow-up out of scope for the low-latency crash. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **Bug Fixes** * Improved low-latency matrix multiplication compatibility across supported GPU architectures. * Prevented execution with unsupported kernel configurations, reducing the risk of invalid tactics and runtime failures. <!-- end of auto-generated comment: release notes by coderabbit.ai -->
PetersonGuo
pushed a commit
to PetersonGuo/flashinfer
that referenced
this pull request
Sep 2, 2026
) (flashinfer-ai#4848) ## 📌 Description Port of flashinfer-ai#4786 (plus the flashinfer-ai#4792 `Sm100f`/sm107 allowance) onto `main`. After flashinfer-ai#4648 the trtllm-gen GEMM pack is a single multi-arch artifact, so `getValidTactics()` on the low-latency runner returned SM107 cubins on Blackwell. Autotune then handed those indices to `cuModuleLoadData`. `trtllm_low_latency_gemm_runner.cu` was the one trtllm-gen runner still missing the `isArchCompatible` / `checkPassingConfigIndex` filter that flashinfer-ai#4280 added to `trtllm_gemm_runner.cu` and `trtllm_batched_gemm_runner.cu`. Do not cherry-pick flashinfer-ai#4786 verbatim: that filter treated `Sm100f` as sm100/sm103 only, and `select_kernel()` still names `_sm100f` heuristics, which flashinfer-ai#4792 showed fails every `mm_fp8` case on Rubin. This PR uses the combined `release-v0.6.18` rule (`Sm100f` on 100/103/107). On B200 (SM100) unpatched `getValidTactics(4, 2560, 8192)` returned 16 indices (`0,2,3,4,5,7,10,11` + `93,95,96,97,101,102,104,109`). Forced tactic `93` failed inside `gemm.run`. After the filter: 8 sm100f tactics; tactic `93` raises `Tactic 93 is not in this runner's compatible config set`. ## 🔍 Related Issues - Closes flashinfer-ai#4773 - Cherry-pick / port of flashinfer-ai#4786 onto `main` (not a verbatim cherry-pick; includes the flashinfer-ai#4792 `Sm100f` sm107 allowance) ## 🚀 Pull Request Checklist ### ✅ Pre-commit Checks - [x] I have installed `pre-commit`. - [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. ## 🧪 Tests - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.). Local B200 (SM100, CUDA 13.0, torch 2.13.0+cu130), worktree at `upstream/main` + this commit: - `pytest tests/gemm/test_mm_fp8.py tests/utils/test_logging_replay.py::test_mm_fp8_replay` → 31 passed - Tactic dump and forced-sm107a FFI path as above Not verified here: SM103 (B300) or SM107 (Rubin). Those were covered on `release-v0.6.18` by flashinfer-ai#4786 / flashinfer-ai#4792. ## Reviewer Notes Sibling runners on `main` still map `Sm100f` to `100 || 103` only. They were left alone: they already have an arch filter, their heuristics have dedicated `_sm107a` names, and flashinfer-ai#4792 called that follow-up out of scope for the low-latency crash. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **Bug Fixes** * Improved low-latency matrix multiplication compatibility across supported GPU architectures. * Prevented execution with unsupported kernel configurations, reducing the risk of invalid tactics and runtime failures. <!-- end of auto-generated comment: release notes by coderabbit.ai -->
5 tasks
1 task
1 task
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
5 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
📌 Description
This PR relands SM 107 support to main branch (reverted in #4171) as well as some other release fixes.
Cherry Picks
Other Changes
TLLM_RUBIN_FEATURES: SiTuGlustatic_asserts + tile-192 advertisement, compiled out for the Rubin BMM pin)
test_unified_moe.pyarch assertions written post-revert(feat(moe): add TRTLLM MXFP4 W4A8 and W4A16 unified API support #4159) flipped to the restored contract (FP4/BF16 claim 107; FP8 stays 100/103)
🔍 Related Issues
#4107, #4164, reverts #4171
🚀 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
pre-commitby runningpip install pre-commit(or used your preferred method).pre-commit install.pre-commit run --all-filesand fixed any reported issues.🧪 Tests
unittest, etc.).Reviewer Notes
Summary by CodeRabbit