Conversation
) The flashinfer-ai#4648 multi-arch GEMM pack lists SM107 configs in the same manifest as SM100. Low-latency autotune loaded those cubins on B200/GB200 and crashed in cuModuleGetFunction. Match the existing trtllm_gemm arch filter so only device-compatible tactics are used.
kahyunnam
requested review from
Aneureka,
aleozlx,
bkryu,
cyx-6,
dhiraj113,
feih-nv,
jiahanc,
jimmyzho,
nv-yunzheq,
qiching,
qsang-nv,
saltyminty,
samuellees,
sricketts,
yongwww,
yyihuang and
yzh119
as code owners
August 27, 2026 19:49
Contributor
|
Important Review skippedAuto reviews are disabled on base/target branches other than the default branch. Please check the settings in the CodeRabbit UI or the ⚙️ Run configurationConfiguration used: defaults Review profile: CHILL Plan: Pro Plus Run ID: You can disable this status message by setting the Use the checkbox below for a quick retry:
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 |
Tactic ids index the multi-arch cubin manifest, so an index that did not come from getValidTactics() (saved autotune config from other hardware, explicit FFI argument) reached cuModuleLoadData and faulted. Mirror trtllm_gemm's checkPassingConfigIndex so it errors instead.
5 tasks
Vinnie6167
added a commit
that referenced
this pull request
Aug 28, 2026
## 📌 Description Five independent SM107 (Rubin) fixes for `release-v0.6.18`, found by a sharded validation sweep of this branch at `89cafe9a4`. Together they remove **133 test failures**. Each commit stands alone and can be reviewed separately. Every fix was A/B-validated on SM107 hardware, comparing **failed, passed and skipped** counts — not just the failure count. ### 1. `fix(moe)` — `check_support()` must not require a bound device Follow-up to #4790, which introduced a regression this branch is carrying. The SM107 CuTe DSL probe read `self.device` **before** deciding whether the architecture was relevant, so `check_support()` raised on a runner with no device attached: ``` AttributeError: 'CuteDslNvfp4Runner' object has no attribute 'device' ``` `tests/moe/test_unified_moe.py::TestMoERunnerSupport` builds runners with `__new__` and attaches only a config — a reasonable way to exercise a pure configuration check. **Not SM107-specific:** the attribute access precedes any compute-capability test, so it raised on every architecture, Blackwell included. Treats a missing device as "nothing arch-specific to decide". `MoELayer` always sets a device in `__init__` before calling `check_support()`, so the dispatch path is unaffected. *Validated: `TestMoERunnerSupport` 64 passed (was 2 failed).* ### 2. `test` — skip SM107 CuTe-DSL cases when the installed DSL predates them Two unrelated failures with one cause: the public stack ships CuTe DSL 4.7.0, which has no `sm_107` in its `Arch` enum and no `cutlass.utils.rubin_helpers`. * **`tests/gemm/test_bmm_fp8.py`** reported a *problem-shape* error for an unavailable backend. `_can_implement_config_sm107` instantiates the kernel class to call `can_implement`; on DSL 4.7.0 that raises `NotImplementedError`, and a surrounding `except Exception: return False` turns it into "this config is invalid" — so every entry of `SM107_AUTOTUNE_CONFIGS` is rejected and the user sees `No valid cute-dsl SM107 bmm_fp8 config for problem (...)`. The geometry was fine. Same pipeline, same test unit, same 7 nodes: internal DSL 4.8 → 7 passed; public DSL 4.7.0 → 6 passed / 1 failed. * **`TestCuteDslMoeW4A16`** is the only GPU-executing DSL class in `tests/moe/test_cute_dsl_fused_moe.py` missing `pytestmark = _requires_dsl_arch` — all 16 classes audited. Its entry point calls `require_cute_dsl_arch(..., native_only=True)`, exactly what that marker tests. *Validated both directions: with 4.7.0 both skip; with 4.8 neither skips and the W4A16 test proceeds into kernel compilation, so it does not over-skip.* ### 3. `test(moe)` — skip the hardcoded unfused-finalize case The SM107 skip fixture keys on parameterization, but `test_deterministic_finalize_numerical_accuracy` passes `use_fused_finalize=False` in its *body*, so it escaped and still hit the `NotImplementedError` the fixture exists to absorb. Matched by function identity, like `test_geglu_tanh_accuracy` — still decided before the body runs, so it cannot absorb a genuine regression. `test_route_tile_boundary_accuracy` and `test_weight_scale_mapping` also hardcode that flag and are **deliberately not skipped**: neither fails in CI, and they use the W4A16 entry point rather than the blockscaled finalize path, so skipping them would drop real coverage. ### 4. `fix(trtllm-gen)` — accept SM107 for `Sm100f` in the low-latency GEMM filter ``` RuntimeError: Check failed: (it != mPassingConfigIndices.end()) is false: Tactic 0 is not in this runner's compatible config set ``` `isArchCompatible` in `csrc/trtllm_low_latency_gemm_runner.cu` mapped `CudaArch::Sm100f` to `smVersion == 100 || smVersion == 103`, omitting 107, so every family-conditional cubin was discarded on Rubin. The mechanism was verified on hardware rather than assumed. `TLLM_RUBIN_FEATURES` **is** defined for the SM107 module (nvcc flags dumped), and the passing set is non-empty unpatched — 8 native `Sm107a` tactics. But `select_kernel()`, twelve lines above the filter, names its heuristic kernels *literally* and every name ends in `_sm100f`, so the index it returns is always one the filter dropped, and `checkPassingConfigIndex()` converts that into the error. Patched, the passing set goes 8 → 16 by adding exactly the `Sm100f` entries. *Validated: `test_mm_fp8` **30 failed → 30 passed**; `test_logging_replay` **1 failed → 16 passed**. Both arms genuinely rebuilt (12 s each, no timing asymmetry). `cos_sim > 0.99` assertions hold on the `Sm100f` kernels.* **Deliberately excluded — the cuDNN mixed-form seqlens failure (64 occurrences).** `_cudnn_supports_direct_seqlens(mixed=True)` authorises a paged path that every currently published `cudnn-frontend` rejects at graph validation, so the request fails with `Padding mask requires seq_len_q/seq_len_kv`. Verified against a wheel installed fresh from PyPI (`nvidia-cudnn-frontend==1.27.0`, the newest published): the mixed form is rejected there too, so this is **not Rubin-specific and not a container artifact** — it affects any architecture taking that path on a stock install. It is excluded from this Rubin-scoped PR and should be fixed separately. A capability-probe fix exists and is validated (32 failed -> 0, passed unchanged) but is held for its own PR. **Deliberately excluded — the PrimTS decode exhaustive-checker race (4 occurrences).** `TmemSoftmaxLocalResource.get_tmem_requirements()` declares a TMEM allocation the kernel never performs when `keeps_stats_via_smem` is set, so the exhaustive checker correctly reports an aliasing race against `tmemS0` and kernel construction fails with `ValueError: Exhaustive checker found 1 aliasing race(s)`. Scope, stated accurately: this is **not** test-only. The production path (`_run_decode_gen_active`) enables the checker via `not (cfg.use_keeps_mma_ab and cfg.num_insts_kv == 1)`, and all four affected profiles measure `num_insts_kv == 2`, so the checker runs for them outside tests too. The failure is a hard construction-time error, not a silent wrong answer, and no runtime behaviour changes either way — a candidate fix was verified to leave TMEM offsets and column counts byte-identical across ten decode profiles. It is excluded here as a deliberate release-management decision: the change is a scheduling-model edit in the PrimTS engine, which is owned elsewhere, and the risk of touching it unfamiliar outweighs a loud, characterised construction failure. A validated minimal backport exists (2 files, +51/-21, 4 failed -> 4 passed) and should be routed to the PrimTS owner rather than landed here. ## 🚀 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.). | Fix | Before → after on SM107 | |---|---| | 1 `check_support` | `TestMoERunnerSupport` 2 failed → **64 passed** | | 2 DSL-4.7 skips | 1 failed → skipped (bmm_fp8); 1 failed → skipped (W4A16) | | 4 unfused-finalize skip | 1 failed → skipped | | 5 low-latency GEMM | 30 failed → **30 passed**; 1 failed → **16 passed** | ## Reviewer Notes * **`Sm100f` is not universally family-valid, and the file says so.** A few lines below the filter changed in commit 5 there is already a carve-out — `sm103` must fall back to `Sm103a` for the f2fp patch. sm107 is empirically fine here (correct numerics, verified), but a similar per-feature gap on Rubin would need the same treatment. * **Two copies of the same defect remain on this branch**, in `csrc/trtllm_gemm_runner.cu` and `csrc/trtllm_batched_gemm_runner.cu`, both still `100 || 103`. They are not producing failures on the current cubin pin and are deliberately out of scope here. * **Commit 5 is release-only by construction.** `csrc/trtllm_low_latency_gemm_runner.cu` has no `isArchCompatible` on `main`; #4773/#4786 added it to this branch only.
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 -->
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
1 participant
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.
Summary
Fixes #4773:
mm_fp8/test_mm_fp8_replaySIGSEGV incuModuleGetFunctionduring 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.hnow lists sm100f and sm107a configs in the manifest the Blackwell module compiles against. Previously the Rubin cubins lived in a separateTRTLLM_GEN_GEMM_RUBINpin, so the non-Rubin manifest was sm100-only.trtllm_low_latency_gemm_runner.cuwas 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);cuModuleLoadDatafails on the first sm107a cubin, theCUresultis ignored by the generatedGemmInterface, andcuModuleGetFunctionfaults on the uninitialisedCUmodule.Two commits, no cubin regeneration:
isArchCompatiblefilter when buildingmPassingConfigIndices, identical tocsrc/trtllm_gemm_runner.cu(Sm107aonly underTLLM_RUBIN_FEATURES,Sm100fallowed on sm100 and sm103).checkPassingConfigIndexinrun(), also mirroringtrtllm_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 viasave_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:
trtllm_gemm_runner.cuisArchCompatible+checkPassingConfigIndex(#4280)trtllm_batched_gemm_runner.cu(trtllm-gen MoE backend)isArchCompatible+checkPassingConfigIndex(#4280)isSMCompatible()infmhaKernels.cuh, with explicit sm107 rulestrtllm_low_latency_gemm_runner.cuTest plan
Local B200 (SM100, CUDA 13.0, Python 3.10), on
release-v0.6.18+ these commits:0,2,3,4,5,7,10,11+ sm107a93,95,96,97,101,102,104,109); SIGSEGV on the first sm107a cubin load.mm_fp8passes underautotune()and on the heuristictactic=-1path.93) now raisesRuntimeError: Tactic 93 is not in this runner's compatible config setinstead of SIGSEGV.pytest tests/gemm/test_mm_fp8.py tests/utils/test_logging_replay.py→ 44 passed, 2 skipped (includestest_mm_fp8_replay, the test that crashed in CI).unit_test_b300/ GB200 jobs coveringtests/gemm/test_mm_fp8.pyandtests/utils/test_logging_replay.py.