Adds SM107 support - #4122
Merged
Merged
Adds SM107 support#4122
Conversation
Contributor
|
Caution The consumer version of Gemini Code Assist on GitHub has been sunset. All code review activity has officially ceased. |
Contributor
|
Note Reviews pausedIt looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the Use the following commands to manage reviews:
Use the checkboxes below for quick actions:
📝 WalkthroughWalkthroughSM107 support is added across CUDA compilation, JIT/AOT generation, backend routing, CuTe DSL checks, GEMM/MoE/attention dispatch, TRTLLM-GEN kernel selection, and architecture-aware tests. TRTLLM attention APIs also gain FP16-softmax and spcompress variant selectors. ChangesSM107 compilation and backend support
TRTLLM-GEN attention variants
Validation and test updates
Estimated code review effort: 5 (Critical) | ~120 minutes Sequence Diagram(s)sequenceDiagram
participant User as FlashInfer API
participant Wrapper as Attention wrapper
participant Launcher as TRTLLM launcher
participant Params as FMHA runner parameters
participant Kernel as FMHA kernel selector
User->>Wrapper: provide use_fp16_softmax / uses_spcompress
Wrapper->>Launcher: forward variant selectors
Launcher->>Params: set mUseFp16Softmax / mUsesSpcompress
Params->>Kernel: include selectors in kernel traits
Kernel-->>Launcher: select matching cubin and launch
Possibly related PRs
Suggested labels: Suggested reviewers: 🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✨ Finishing Touches🧪 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
force-pushed
the
sm107-support
branch
from
July 24, 2026 02:20
e52db4d to
26b5902
Compare
Contributor
|
/bot run |
Collaborator
|
GitLab MR !1032 has been created, and the CI pipeline #59351310 is currently running. I'll report back once the pipeline job completes. |
jimmyzho
marked this pull request as ready for review
July 24, 2026 06:30
Contributor
|
Caution The consumer version of Gemini Code Assist on GitHub has been sunset. All code review activity has officially ceased. |
jimmyzho
requested review from
IwakuraRein,
aleozlx,
bkryu,
cyx-6,
dhiraj113,
jiahanc,
jimmyzho,
kahyunnam,
nv-yunzheq,
qsang-nv,
saltyminty,
samuellees,
sricketts,
yongwww,
yyihuang and
yzh119
as code owners
July 24, 2026 06:31
5 tasks
aleozlx
added a commit
that referenced
this pull request
Jul 28, 2026
…ild break (#4200) <!-- .github/pull_request_template.md --> ## 📌 Description 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 #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 #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, #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. ## 🔍 Related Issues - Release run: https://github.com/flashinfer-ai/flashinfer/actions/runs/30330417436 - Introduced by the per-arch pin split (#4191) interacting with SM107 support (#4122) ## 🚀 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.). `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 ``` ## Reviewer Notes - **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>
This was referenced Jul 28, 2026
kahyunnam
pushed a commit
that referenced
this pull request
Jul 28, 2026
## 📌 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 #4168 / #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 #4122.
🤖 Generated with [Claude Code](https://claude.com/claude-code)
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
kahyunnam
added a commit
that referenced
this pull request
Jul 29, 2026
…(release port for #4107) (#4230) ## Description Release-specific port of #4177 onto `release-v0.6.16` for #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`. ## Changes - **`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. ## Release-specific notes - Includes **sm107** in the allowlist (unlike the rebased #4177 for `main`, which will drop 107 after the #4122 revert). - `Sm107a` enum case is gated behind `TLLM_RUBIN_FEATURES`, matching the existing release pattern from #4122/#4191. Verified against the actual pinned headers: the default BMM/GEMM pins do not define `Sm107a`; only the Rubin pins do. ## Verification - 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 ## Related - Fixes #4107 - Upstream PR: #4177 (targets `main`, has merge conflicts) - Cherry-pick source commit on `main`: not yet merged ## Pre-existing issue (not in scope) `#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.
5 tasks
1 task
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>
5 tasks
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>
5 tasks
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>
jimmyzho
added a commit
that referenced
this pull request
Aug 20, 2026
…4596) <!-- .github/pull_request_template.md --> ## 📌 Description This adds those two attributes to the kernel hash and restores the selection knobs that go with them, so the variants are both distinguishable and addressable. It re-lands the trtllm-gen FMHA changes reverted in c04ed08 ("Revert trtllm-gen FMHA rubin-specific features", #4122), rebased onto the current launcher. ## Changes - **Hash** (`fmhaKernels.cuh`): hash `mFp16Softmax` and `mUsesSpcompress`. The key was full at bit 62, so the enum fields are now packed to their real width (`qkvLayout` 2, `maskType` 3, `kernelType` 3, `tileScheduler` 2) instead of a round 4 bits each. That frees the two bits and leaves 57–63 spare. The tightened fields are range-checked so an out-of-range enum fails loudly rather than aliasing onto its neighbour. The key is a private lookup key that nothing persists, so the layout is free to move as long as both hash entry points move with it. - **Params** (`fmhaRunnerParams.h`): `mUseFp16Softmax` / `mUsesSpcompress` on `TllmGenFmhaRunnerParams` and `TllmGenSelectKernelParams`. - **Plumbing** (`trtllm_fmha_kernel_launcher.cu`, `prefill.py`, `mla/_core.py`, `decode.py`): `use_fp16_softmax` / `uses_spcompress` through the paged/ragged/context launchers, the prefill wrappers, and MLA decode — including the cute-dsl incompatibility check that rejects `use_fp16_softmax` on that backend. - **Tests** (`test_trtllm_gen_mla.py`): `use_fp16_softmax` coverage from the original change. ## 🔍 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 - **New Features** - Added optional FP16 softmax selection for supported TRT-LLM paged, ragged, context, decode, and MLA attention workflows. - Added optional sparse-compression kernel selection for supported paged and ragged attention workflows. - Options remain disabled by default and are unavailable on unsupported backends or hardware. - **Bug Fixes** - Improved attention kernel selection across supported configurations. - Added clear validation errors when SM107-only features are used on incompatible devices. - **Tests** - Expanded coverage across model dimensions, batch sizes, page sizes, and query lengths. <!-- end of auto-generated comment: release notes by coderabbit.ai -->
kahyunnam
pushed a commit
that referenced
this pull request
Aug 20, 2026
…4596) <!-- .github/pull_request_template.md --> ## 📌 Description This adds those two attributes to the kernel hash and restores the selection knobs that go with them, so the variants are both distinguishable and addressable. It re-lands the trtllm-gen FMHA changes reverted in c04ed08 ("Revert trtllm-gen FMHA rubin-specific features", #4122), rebased onto the current launcher. ## Changes - **Hash** (`fmhaKernels.cuh`): hash `mFp16Softmax` and `mUsesSpcompress`. The key was full at bit 62, so the enum fields are now packed to their real width (`qkvLayout` 2, `maskType` 3, `kernelType` 3, `tileScheduler` 2) instead of a round 4 bits each. That frees the two bits and leaves 57–63 spare. The tightened fields are range-checked so an out-of-range enum fails loudly rather than aliasing onto its neighbour. The key is a private lookup key that nothing persists, so the layout is free to move as long as both hash entry points move with it. - **Params** (`fmhaRunnerParams.h`): `mUseFp16Softmax` / `mUsesSpcompress` on `TllmGenFmhaRunnerParams` and `TllmGenSelectKernelParams`. - **Plumbing** (`trtllm_fmha_kernel_launcher.cu`, `prefill.py`, `mla/_core.py`, `decode.py`): `use_fp16_softmax` / `uses_spcompress` through the paged/ragged/context launchers, the prefill wrappers, and MLA decode — including the cute-dsl incompatibility check that rejects `use_fp16_softmax` on that backend. - **Tests** (`test_trtllm_gen_mla.py`): `use_fp16_softmax` coverage from the original change. ## 🔍 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 - **New Features** - Added optional FP16 softmax selection for supported TRT-LLM paged, ragged, context, decode, and MLA attention workflows. - Added optional sparse-compression kernel selection for supported paged and ragged attention workflows. - Options remain disabled by default and are unavailable on unsupported backends or hardware. - **Bug Fixes** - Improved attention kernel selection across supported configurations. - Added clear validation errors when SM107-only features are used on incompatible devices. - **Tests** - Expanded coverage across model dimensions, batch sizes, page sizes, and query lengths. <!-- end of auto-generated comment: release notes by coderabbit.ai --> (cherry picked from commit ad8bb37)
kahyunnam
pushed a commit
that referenced
this pull request
Aug 21, 2026
…4596) <!-- .github/pull_request_template.md --> ## 📌 Description This adds those two attributes to the kernel hash and restores the selection knobs that go with them, so the variants are both distinguishable and addressable. It re-lands the trtllm-gen FMHA changes reverted in c04ed08 ("Revert trtllm-gen FMHA rubin-specific features", #4122), rebased onto the current launcher. ## Changes - **Hash** (`fmhaKernels.cuh`): hash `mFp16Softmax` and `mUsesSpcompress`. The key was full at bit 62, so the enum fields are now packed to their real width (`qkvLayout` 2, `maskType` 3, `kernelType` 3, `tileScheduler` 2) instead of a round 4 bits each. That frees the two bits and leaves 57–63 spare. The tightened fields are range-checked so an out-of-range enum fails loudly rather than aliasing onto its neighbour. The key is a private lookup key that nothing persists, so the layout is free to move as long as both hash entry points move with it. - **Params** (`fmhaRunnerParams.h`): `mUseFp16Softmax` / `mUsesSpcompress` on `TllmGenFmhaRunnerParams` and `TllmGenSelectKernelParams`. - **Plumbing** (`trtllm_fmha_kernel_launcher.cu`, `prefill.py`, `mla/_core.py`, `decode.py`): `use_fp16_softmax` / `uses_spcompress` through the paged/ragged/context launchers, the prefill wrappers, and MLA decode — including the cute-dsl incompatibility check that rejects `use_fp16_softmax` on that backend. - **Tests** (`test_trtllm_gen_mla.py`): `use_fp16_softmax` coverage from the original change. ## 🔍 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 - **New Features** - Added optional FP16 softmax selection for supported TRT-LLM paged, ragged, context, decode, and MLA attention workflows. - Added optional sparse-compression kernel selection for supported paged and ragged attention workflows. - Options remain disabled by default and are unavailable on unsupported backends or hardware. - **Bug Fixes** - Improved attention kernel selection across supported configurations. - Added clear validation errors when SM107-only features are used on incompatible devices. - **Tests** - Expanded coverage across model dimensions, batch sizes, page sizes, and query lengths. <!-- end of auto-generated comment: release notes by coderabbit.ai --> (cherry picked from commit ad8bb37)
Vinnie6167
added a commit
to Vinnie6167/flashinfer
that referenced
this pull request
Aug 21, 2026
supported_compute_capability gates on hardware capability alone, so flashinfer-ai#4122 widening the cute-dsl lists to [100, 103, 107] told the dispatcher that Rubin is supported regardless of which CuTe DSL is installed. On DSL 4.7.0 -- which has no sm_107a in its Arch enum -- that produces `KeyError: 'sm_107a'` from enum.py inside cute.compile, with no flashinfer frame in the traceback and no warning. Two auto heuristics route sm_107 toward cute-dsl specifically, so this is reachable without asking for it: _heuristic_func_mm_fp4 has an `elif is_sm107: candidate_backends = (..., "cute-dsl")` branch, and _heuristic_func_bmm_fp8 appends "cute-dsl_sm107" when is_sm107_supported. Make 107 conditional for the three cute-dsl backends that claim it. The predicate is is_cute_dsl_arch_supported(10, 7), which is already exactly the intended rule: true on DSL 4.8+, or on older DSL when the user exported CUTE_DSL_ARCH=sm_100f before the process started. FlashInfer does not set that variable itself -- it is the user's choice, since it retargets every DSL kernel in the process. supported_compute_capability grows an optional conditional_ccs mapping. It must be lazy: probing the DSL imports cutlass, which cannot become a hard dependency of `import flashinfer`, and the decorator runs at import time. _supported_ccs stays a plain set when nothing is conditional, so existing comparisons and iteration (tests/utils/test_decorators.py, tests/grouped_mm/conftest.py) are unaffected. A predicate that raises counts as unsupported: advertising a capability we could not verify is what crashes inside kernel compilation. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
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>
Vinnie6167
added a commit
to Vinnie6167/flashinfer
that referenced
this pull request
Aug 21, 2026
supported_compute_capability gates on hardware capability alone, so flashinfer-ai#4122 widening the cute-dsl lists to [100, 103, 107] told the dispatcher Rubin is supported regardless of which CuTe DSL is installed. On DSL 4.7.0 -- no sm_107a in its Arch enum -- that yields `KeyError: 'sm_107a'` from enum.py inside cute.compile, with no flashinfer frame in the traceback and no warning. Two auto heuristics route sm_107 toward cute-dsl specifically, so this is reachable without asking for it: _heuristic_func_mm_fp4 has an `elif is_sm107: candidate_backends = (..., "cute-dsl")` branch, and _heuristic_func_bmm_fp8 appends "cute-dsl_sm107" when is_sm107_supported. Make 107 conditional for the three cute-dsl backends that claim it. The predicate is is_cute_dsl_arch_supported(10, 7) -- exactly the intended rule: true on DSL 4.8+, or on older DSL when the user exported CUTE_DSL_ARCH=sm_100f before the process started. supported_compute_capability grows an optional conditional_ccs mapping. It must be lazy: probing the DSL imports cutlass, which cannot become a hard dependency of `import flashinfer`, and the decorator runs at import time. _supported_ccs stays a plain set when nothing is conditional, so existing comparisons and iteration (tests/utils/test_decorators.py, tests/grouped_mm/conftest.py) are unaffected. A predicate that raises counts as unsupported: advertising a capability we could not verify is what crashes inside kernel compilation. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Vinnie6167
added a commit
to Vinnie6167/flashinfer
that referenced
this pull request
Aug 21, 2026
_check_cute_dsl_arch had reimplemented what flashinfer-ai#4122 already added: require_cute_dsl_arch(device) does the same device -> capability lookup and predicate call, and builds a better message -- it derives the family arch and names the exact value to export, e.g. "CUTE_DSL_ARCH=sm_100f", rather than the generic hint I had written. Delegate to it and adapt only the exception type. That part is load bearing and is why the call cannot be used unchanged: it raises NotImplementedError, while suitable_auto_backends catches ValueError to mean "backend not suitable". Left as-is, an unsupported DSL would propagate out of the auto path and fail the call instead of falling back to cutlass/cudnn. Verified the adapter: passes through when supported, converts NotImplementedError to ValueError with the message intact, and stays silent when the probe itself cannot be imported. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
bkryu
pushed a commit
that referenced
this pull request
Aug 27, 2026
…4649) ## Problem On SM107 (Rubin) with a CuTe DSL older than 4.8, FlashInfer fails with a bare `KeyError: 'sm_107a'` raised from `enum.py` inside `cute.compile` — no FlashInfer frame in the traceback, and no warning. The root cause is that `supported_compute_capability` gates on **hardware capability alone**. #4122 widened the cute-dsl lists to `[100, 103, 107]`, which tells the dispatcher Rubin is supported regardless of which DSL is installed. Public CuTe DSL tops out at 4.7.0 on PyPI, and that release has no `sm_107a` in its `Arch` enum. This is reachable without asking for cute-dsl explicitly — two `auto` heuristics route SM107 *toward* it: - `_heuristic_func_mm_fp4`: `elif is_sm107: candidate_backends = ("cudnn", "cutlass", "cute-dsl")` - `_heuristic_func_bmm_fp8`: appends `"cute-dsl_sm107"` when `is_sm107_supported` ## Changes **1. Decline the cute-dsl backend when the installed DSL cannot emit for the device** (`fix(gemm)`) The three cute-dsl requirement functions now call `_check_cute_dsl_arch(...)`, which sits beside the existing `_check_cute_dsl_availability()` and delegates to `require_cute_dsl_arch()` — the helper added in #4122, which owns both the predicate and the message (it derives the family arch and names the exact `CUTE_DSL_ARCH=sm_100f` to export). Only the exception type is adapted, and that part is load bearing: `require_cute_dsl_arch` raises `NotImplementedError`, while `suitable_auto_backends` catches `ValueError` to mean "backend not suitable" and keeps searching. Left unadapted, an unsupported DSL would propagate out of the auto path and fail the call instead of falling back to cutlass/cudnn. Returning `False` instead of raising was also rejected: on the explicit-backend path that surfaces as `ValueError: Problem size is not supported`, which is misleading. **No capability lists change.** This is deliberately an *availability* check, not a capability one. The kernels do exist for sm_107, so `@supported_compute_capability([100, 103, 107])` stays as-is and `is_backend_supported("cute-dsl", 107)` keeps answering `True` — it is a public method on the wrapper, called with no tensors by e.g. `flashinfer/trace/templates/gemm.py:707`, and making it vary with an installed pip package would also have made the skip reason in `tests/grouped_mm/conftest.py` environment-dependent. This mirrors how the codebase already separates the two axes: `_cudnn_mm_mxfp8_requirement` lists its capabilities statically while `CUDNN_AVAILABLE` handles presence, and `_is_cudnn_override_shape_available` handles a dependency that is present but too old. **2. GDN CP delta rule resolves the arch instead of formatting it** (`fix(gdn)`) `_blackwell_compile_options` guards on the major only, then builds `f"sm_{major}{minor}a"`. Rubin is 10.7, so it passes a check written when "compute 10.x" meant Blackwell 10.0/10.3. This is the only place FlashInfer names the arch for a compute-10 device; everywhere else the DSL derives it internally. `cute_dsl_compile_arch()` returns the device's own arch when the DSL has it, the family arch when the DSL is targeting `sm_100f`, and otherwise raises `NotImplementedError` naming `CUTE_DSL_ARCH`. Same rule as the capability gate, so the two cannot disagree. ## Testing Rubin CI, `TEST_PATH="tests/gemm tests/gdn"`, against `release-v0.6.18`, with `CUTE_DSL_ARCH=sm_100f` exported and public CuTe DSL 4.7.0: | | before | after | |---|---|---| | passed | 8,141 | **12,341** | | failed | 4,230 | **1** | | `KeyError: 'sm_107a'` | 4,482 | **0** | Identical results on **both** VR200 (`hecate`, 4 workers, 2,078s) and GR100 (8 workers, 3,424s); `suite_complete=true` on both, well inside the 13,500s deadline. Per-file, verified independently on both boards: | File | before | after | |---|---|---| | `tests/gdn/test_prefill_delta_rule.py` | 2,678 | **0** | | `tests/gemm/test_mm_mxfp8.py` | 501 | **0** | | `tests/gdn/test_decode_delta_rule.py` | 417 | **0** | | `tests/gdn/test_prefill_cp_delta_rule.py` | 232 | **0** | The remaining failures are `tests/gdn/test_decode_ucache.py` and `tests/gemm/test_bmm_fp8.py` — see below. `BackendSupportedError` count is **0**, so the cute-dsl backends are being selected and compiling successfully against `sm_100f` — not silently skipped. The node accounting reconciles exactly: the plan drops 44,275 → 44,229 nodes and 37 → 36 units, i.e. the 46 tests in the ucache module leave collection entirely (a module-level skip is taken during collection, so those nodes are not counted as `skipped`). `passed` moves +287 = +288 Triton tests now compiling, −1 ucache test that previously passed; `failed` moves −333 = −288 Triton −45 ucache. Also unit-tested away from hardware: the new decorator resolves conditional 107 as False on DSL 4.7, True on 4.8+/`CUTE_DSL_ARCH`, False when the predicate raises, and yields a plain `set` when no conditional is given. `cute_dsl_compile_arch` was verified against a stubbed `Arch` enum for native / family / unsupported / Blackwell-unchanged, and the skip predicate for all four DSL-vs-arch combinations. ### Caveats - **The numbers above do not reflect this branch.** They were measured at `93143db2`, which carried a skip guard for `tests/gdn/test_decode_ucache.py` that has since been reverted, so 45 of those tests now fail again rather than skipping. - **The gemm mechanism changed after that measurement.** The two commits after it moved the check out of the decorator and into the requirement functions; that mechanism is unit-tested (adapter pass-through, `NotImplementedError` → `ValueError`, silent when the probe cannot be imported) but has not been re-run on hardware. - The measurement runs also carry `CUTE_DSL_ARCH=sm_100f` from the CI side. With it set the DSL *can* target sm_107, so `_check_cute_dsl_arch` passes and the gemm change is a no-op; only a run without that variable exercises the deselect-and-fall-back path. - `cute_dsl_compile_arch` changes `gdn_cp_prefill.py` for **all** compute-10 devices, not just Rubin. Blackwell resolution (`sm_100a` / `sm_103a`) is verified against a stubbed `Arch` enum, not on B200/GB200 hardware. ## Not addressed - **45 `KeyError: 'sm_107a'`** in `tests/gdn/test_decode_ucache.py`. Not fixable from FlashInfer: those kernels compile through `@cute.experimental.jit` / `@cute.experimental.kernel`, passing no arch and no compile options, so the DSL resolves the device arch itself and looks up `sm_107a` in its own enum. There is no FlashInfer-side site to guard, the traceback bottoms out at `enum.py:813` with no FlashInfer frame, and `CUTE_DSL_ARCH=sm_100f` does not help because that path never consults it — which points at a genuine **CuTe DSL 4.8** requirement. Left visible rather than skipped; the kernel author (#4081) is better placed to say whether it is inherent. - **1 `No valid cute-dsl SM107 bmm_fp8 config`** in `tests/gemm/test_bmm_fp8.py` — pre-existing, and present on the internal DSL 4.8 stack too (18 vs 20 occurrences across stacks), so it is independent of the DSL version question. The 288 Triton `PTXASError` failures previously seen in `tests/gemm/test_group_gemm.py` were a CI-side issue, not a FlashInfer one: Triton resolves ptxas through its own knobs (`TRITON_PTXAS_PATH`, and `TRITON_PTXAS_BLACKWELL_PATH` for arch >= 100, which is the one Rubin selects) and otherwise falls back to `$CUDA_HOME/bin/ptxas`. Fixed in flashinfer-ci!354; this run confirms 0 remaining. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **New Features** * Added architecture detection for CuTe DSL compilation, including native and family-compatible GPU architectures. * Added clear guidance when the installed DSL cannot compile for a target GPU. * **Bug Fixes** * Improved Blackwell architecture handling, including support for devices with nonstandard architecture identifiers. * Prevented unsuitable CuTe DSL backends from being selected automatically when architecture support is unavailable. <!-- end of auto-generated comment: release notes by coderabbit.ai --> Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
PetersonGuo
pushed a commit
to PetersonGuo/flashinfer
that referenced
this pull request
Aug 28, 2026
…lashinfer-ai#4649) ## Problem On SM107 (Rubin) with a CuTe DSL older than 4.8, FlashInfer fails with a bare `KeyError: 'sm_107a'` raised from `enum.py` inside `cute.compile` — no FlashInfer frame in the traceback, and no warning. The root cause is that `supported_compute_capability` gates on **hardware capability alone**. flashinfer-ai#4122 widened the cute-dsl lists to `[100, 103, 107]`, which tells the dispatcher Rubin is supported regardless of which DSL is installed. Public CuTe DSL tops out at 4.7.0 on PyPI, and that release has no `sm_107a` in its `Arch` enum. This is reachable without asking for cute-dsl explicitly — two `auto` heuristics route SM107 *toward* it: - `_heuristic_func_mm_fp4`: `elif is_sm107: candidate_backends = ("cudnn", "cutlass", "cute-dsl")` - `_heuristic_func_bmm_fp8`: appends `"cute-dsl_sm107"` when `is_sm107_supported` ## Changes **1. Decline the cute-dsl backend when the installed DSL cannot emit for the device** (`fix(gemm)`) The three cute-dsl requirement functions now call `_check_cute_dsl_arch(...)`, which sits beside the existing `_check_cute_dsl_availability()` and delegates to `require_cute_dsl_arch()` — the helper added in flashinfer-ai#4122, which owns both the predicate and the message (it derives the family arch and names the exact `CUTE_DSL_ARCH=sm_100f` to export). Only the exception type is adapted, and that part is load bearing: `require_cute_dsl_arch` raises `NotImplementedError`, while `suitable_auto_backends` catches `ValueError` to mean "backend not suitable" and keeps searching. Left unadapted, an unsupported DSL would propagate out of the auto path and fail the call instead of falling back to cutlass/cudnn. Returning `False` instead of raising was also rejected: on the explicit-backend path that surfaces as `ValueError: Problem size is not supported`, which is misleading. **No capability lists change.** This is deliberately an *availability* check, not a capability one. The kernels do exist for sm_107, so `@supported_compute_capability([100, 103, 107])` stays as-is and `is_backend_supported("cute-dsl", 107)` keeps answering `True` — it is a public method on the wrapper, called with no tensors by e.g. `flashinfer/trace/templates/gemm.py:707`, and making it vary with an installed pip package would also have made the skip reason in `tests/grouped_mm/conftest.py` environment-dependent. This mirrors how the codebase already separates the two axes: `_cudnn_mm_mxfp8_requirement` lists its capabilities statically while `CUDNN_AVAILABLE` handles presence, and `_is_cudnn_override_shape_available` handles a dependency that is present but too old. **2. GDN CP delta rule resolves the arch instead of formatting it** (`fix(gdn)`) `_blackwell_compile_options` guards on the major only, then builds `f"sm_{major}{minor}a"`. Rubin is 10.7, so it passes a check written when "compute 10.x" meant Blackwell 10.0/10.3. This is the only place FlashInfer names the arch for a compute-10 device; everywhere else the DSL derives it internally. `cute_dsl_compile_arch()` returns the device's own arch when the DSL has it, the family arch when the DSL is targeting `sm_100f`, and otherwise raises `NotImplementedError` naming `CUTE_DSL_ARCH`. Same rule as the capability gate, so the two cannot disagree. ## Testing Rubin CI, `TEST_PATH="tests/gemm tests/gdn"`, against `release-v0.6.18`, with `CUTE_DSL_ARCH=sm_100f` exported and public CuTe DSL 4.7.0: | | before | after | |---|---|---| | passed | 8,141 | **12,341** | | failed | 4,230 | **1** | | `KeyError: 'sm_107a'` | 4,482 | **0** | Identical results on **both** VR200 (`hecate`, 4 workers, 2,078s) and GR100 (8 workers, 3,424s); `suite_complete=true` on both, well inside the 13,500s deadline. Per-file, verified independently on both boards: | File | before | after | |---|---|---| | `tests/gdn/test_prefill_delta_rule.py` | 2,678 | **0** | | `tests/gemm/test_mm_mxfp8.py` | 501 | **0** | | `tests/gdn/test_decode_delta_rule.py` | 417 | **0** | | `tests/gdn/test_prefill_cp_delta_rule.py` | 232 | **0** | The remaining failures are `tests/gdn/test_decode_ucache.py` and `tests/gemm/test_bmm_fp8.py` — see below. `BackendSupportedError` count is **0**, so the cute-dsl backends are being selected and compiling successfully against `sm_100f` — not silently skipped. The node accounting reconciles exactly: the plan drops 44,275 → 44,229 nodes and 37 → 36 units, i.e. the 46 tests in the ucache module leave collection entirely (a module-level skip is taken during collection, so those nodes are not counted as `skipped`). `passed` moves +287 = +288 Triton tests now compiling, −1 ucache test that previously passed; `failed` moves −333 = −288 Triton −45 ucache. Also unit-tested away from hardware: the new decorator resolves conditional 107 as False on DSL 4.7, True on 4.8+/`CUTE_DSL_ARCH`, False when the predicate raises, and yields a plain `set` when no conditional is given. `cute_dsl_compile_arch` was verified against a stubbed `Arch` enum for native / family / unsupported / Blackwell-unchanged, and the skip predicate for all four DSL-vs-arch combinations. ### Caveats - **The numbers above do not reflect this branch.** They were measured at `93143db2`, which carried a skip guard for `tests/gdn/test_decode_ucache.py` that has since been reverted, so 45 of those tests now fail again rather than skipping. - **The gemm mechanism changed after that measurement.** The two commits after it moved the check out of the decorator and into the requirement functions; that mechanism is unit-tested (adapter pass-through, `NotImplementedError` → `ValueError`, silent when the probe cannot be imported) but has not been re-run on hardware. - The measurement runs also carry `CUTE_DSL_ARCH=sm_100f` from the CI side. With it set the DSL *can* target sm_107, so `_check_cute_dsl_arch` passes and the gemm change is a no-op; only a run without that variable exercises the deselect-and-fall-back path. - `cute_dsl_compile_arch` changes `gdn_cp_prefill.py` for **all** compute-10 devices, not just Rubin. Blackwell resolution (`sm_100a` / `sm_103a`) is verified against a stubbed `Arch` enum, not on B200/GB200 hardware. ## Not addressed - **45 `KeyError: 'sm_107a'`** in `tests/gdn/test_decode_ucache.py`. Not fixable from FlashInfer: those kernels compile through `@cute.experimental.jit` / `@cute.experimental.kernel`, passing no arch and no compile options, so the DSL resolves the device arch itself and looks up `sm_107a` in its own enum. There is no FlashInfer-side site to guard, the traceback bottoms out at `enum.py:813` with no FlashInfer frame, and `CUTE_DSL_ARCH=sm_100f` does not help because that path never consults it — which points at a genuine **CuTe DSL 4.8** requirement. Left visible rather than skipped; the kernel author (flashinfer-ai#4081) is better placed to say whether it is inherent. - **1 `No valid cute-dsl SM107 bmm_fp8 config`** in `tests/gemm/test_bmm_fp8.py` — pre-existing, and present on the internal DSL 4.8 stack too (18 vs 20 occurrences across stacks), so it is independent of the DSL version question. The 288 Triton `PTXASError` failures previously seen in `tests/gemm/test_group_gemm.py` were a CI-side issue, not a FlashInfer one: Triton resolves ptxas through its own knobs (`TRITON_PTXAS_PATH`, and `TRITON_PTXAS_BLACKWELL_PATH` for arch >= 100, which is the one Rubin selects) and otherwise falls back to `$CUDA_HOME/bin/ptxas`. Fixed in flashinfer-ci!354; this run confirms 0 remaining. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **New Features** * Added architecture detection for CuTe DSL compilation, including native and family-compatible GPU architectures. * Added clear guidance when the installed DSL cannot compile for a target GPU. * **Bug Fixes** * Improved Blackwell architecture handling, including support for devices with nonstandard architecture identifiers. * Prevented unsuitable CuTe DSL backends from being selected automatically when architecture support is unavailable. <!-- end of auto-generated comment: release notes by coderabbit.ai --> Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
PetersonGuo
pushed a commit
to PetersonGuo/flashinfer
that referenced
this pull request
Aug 28, 2026
…lashinfer-ai#4649) ## Problem On SM107 (Rubin) with a CuTe DSL older than 4.8, FlashInfer fails with a bare `KeyError: 'sm_107a'` raised from `enum.py` inside `cute.compile` — no FlashInfer frame in the traceback, and no warning. The root cause is that `supported_compute_capability` gates on **hardware capability alone**. flashinfer-ai#4122 widened the cute-dsl lists to `[100, 103, 107]`, which tells the dispatcher Rubin is supported regardless of which DSL is installed. Public CuTe DSL tops out at 4.7.0 on PyPI, and that release has no `sm_107a` in its `Arch` enum. This is reachable without asking for cute-dsl explicitly — two `auto` heuristics route SM107 *toward* it: - `_heuristic_func_mm_fp4`: `elif is_sm107: candidate_backends = ("cudnn", "cutlass", "cute-dsl")` - `_heuristic_func_bmm_fp8`: appends `"cute-dsl_sm107"` when `is_sm107_supported` ## Changes **1. Decline the cute-dsl backend when the installed DSL cannot emit for the device** (`fix(gemm)`) The three cute-dsl requirement functions now call `_check_cute_dsl_arch(...)`, which sits beside the existing `_check_cute_dsl_availability()` and delegates to `require_cute_dsl_arch()` — the helper added in flashinfer-ai#4122, which owns both the predicate and the message (it derives the family arch and names the exact `CUTE_DSL_ARCH=sm_100f` to export). Only the exception type is adapted, and that part is load bearing: `require_cute_dsl_arch` raises `NotImplementedError`, while `suitable_auto_backends` catches `ValueError` to mean "backend not suitable" and keeps searching. Left unadapted, an unsupported DSL would propagate out of the auto path and fail the call instead of falling back to cutlass/cudnn. Returning `False` instead of raising was also rejected: on the explicit-backend path that surfaces as `ValueError: Problem size is not supported`, which is misleading. **No capability lists change.** This is deliberately an *availability* check, not a capability one. The kernels do exist for sm_107, so `@supported_compute_capability([100, 103, 107])` stays as-is and `is_backend_supported("cute-dsl", 107)` keeps answering `True` — it is a public method on the wrapper, called with no tensors by e.g. `flashinfer/trace/templates/gemm.py:707`, and making it vary with an installed pip package would also have made the skip reason in `tests/grouped_mm/conftest.py` environment-dependent. This mirrors how the codebase already separates the two axes: `_cudnn_mm_mxfp8_requirement` lists its capabilities statically while `CUDNN_AVAILABLE` handles presence, and `_is_cudnn_override_shape_available` handles a dependency that is present but too old. **2. GDN CP delta rule resolves the arch instead of formatting it** (`fix(gdn)`) `_blackwell_compile_options` guards on the major only, then builds `f"sm_{major}{minor}a"`. Rubin is 10.7, so it passes a check written when "compute 10.x" meant Blackwell 10.0/10.3. This is the only place FlashInfer names the arch for a compute-10 device; everywhere else the DSL derives it internally. `cute_dsl_compile_arch()` returns the device's own arch when the DSL has it, the family arch when the DSL is targeting `sm_100f`, and otherwise raises `NotImplementedError` naming `CUTE_DSL_ARCH`. Same rule as the capability gate, so the two cannot disagree. ## Testing Rubin CI, `TEST_PATH="tests/gemm tests/gdn"`, against `release-v0.6.18`, with `CUTE_DSL_ARCH=sm_100f` exported and public CuTe DSL 4.7.0: | | before | after | |---|---|---| | passed | 8,141 | **12,341** | | failed | 4,230 | **1** | | `KeyError: 'sm_107a'` | 4,482 | **0** | Identical results on **both** VR200 (`hecate`, 4 workers, 2,078s) and GR100 (8 workers, 3,424s); `suite_complete=true` on both, well inside the 13,500s deadline. Per-file, verified independently on both boards: | File | before | after | |---|---|---| | `tests/gdn/test_prefill_delta_rule.py` | 2,678 | **0** | | `tests/gemm/test_mm_mxfp8.py` | 501 | **0** | | `tests/gdn/test_decode_delta_rule.py` | 417 | **0** | | `tests/gdn/test_prefill_cp_delta_rule.py` | 232 | **0** | The remaining failures are `tests/gdn/test_decode_ucache.py` and `tests/gemm/test_bmm_fp8.py` — see below. `BackendSupportedError` count is **0**, so the cute-dsl backends are being selected and compiling successfully against `sm_100f` — not silently skipped. The node accounting reconciles exactly: the plan drops 44,275 → 44,229 nodes and 37 → 36 units, i.e. the 46 tests in the ucache module leave collection entirely (a module-level skip is taken during collection, so those nodes are not counted as `skipped`). `passed` moves +287 = +288 Triton tests now compiling, −1 ucache test that previously passed; `failed` moves −333 = −288 Triton −45 ucache. Also unit-tested away from hardware: the new decorator resolves conditional 107 as False on DSL 4.7, True on 4.8+/`CUTE_DSL_ARCH`, False when the predicate raises, and yields a plain `set` when no conditional is given. `cute_dsl_compile_arch` was verified against a stubbed `Arch` enum for native / family / unsupported / Blackwell-unchanged, and the skip predicate for all four DSL-vs-arch combinations. ### Caveats - **The numbers above do not reflect this branch.** They were measured at `93143db2`, which carried a skip guard for `tests/gdn/test_decode_ucache.py` that has since been reverted, so 45 of those tests now fail again rather than skipping. - **The gemm mechanism changed after that measurement.** The two commits after it moved the check out of the decorator and into the requirement functions; that mechanism is unit-tested (adapter pass-through, `NotImplementedError` → `ValueError`, silent when the probe cannot be imported) but has not been re-run on hardware. - The measurement runs also carry `CUTE_DSL_ARCH=sm_100f` from the CI side. With it set the DSL *can* target sm_107, so `_check_cute_dsl_arch` passes and the gemm change is a no-op; only a run without that variable exercises the deselect-and-fall-back path. - `cute_dsl_compile_arch` changes `gdn_cp_prefill.py` for **all** compute-10 devices, not just Rubin. Blackwell resolution (`sm_100a` / `sm_103a`) is verified against a stubbed `Arch` enum, not on B200/GB200 hardware. ## Not addressed - **45 `KeyError: 'sm_107a'`** in `tests/gdn/test_decode_ucache.py`. Not fixable from FlashInfer: those kernels compile through `@cute.experimental.jit` / `@cute.experimental.kernel`, passing no arch and no compile options, so the DSL resolves the device arch itself and looks up `sm_107a` in its own enum. There is no FlashInfer-side site to guard, the traceback bottoms out at `enum.py:813` with no FlashInfer frame, and `CUTE_DSL_ARCH=sm_100f` does not help because that path never consults it — which points at a genuine **CuTe DSL 4.8** requirement. Left visible rather than skipped; the kernel author (flashinfer-ai#4081) is better placed to say whether it is inherent. - **1 `No valid cute-dsl SM107 bmm_fp8 config`** in `tests/gemm/test_bmm_fp8.py` — pre-existing, and present on the internal DSL 4.8 stack too (18 vs 20 occurrences across stacks), so it is independent of the DSL version question. The 288 Triton `PTXASError` failures previously seen in `tests/gemm/test_group_gemm.py` were a CI-side issue, not a FlashInfer one: Triton resolves ptxas through its own knobs (`TRITON_PTXAS_PATH`, and `TRITON_PTXAS_BLACKWELL_PATH` for arch >= 100, which is the one Rubin selects) and otherwise falls back to `$CUDA_HOME/bin/ptxas`. Fixed in flashinfer-ci!354; this run confirms 0 remaining. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **New Features** * Added architecture detection for CuTe DSL compilation, including native and family-compatible GPU architectures. * Added clear guidance when the installed DSL cannot compile for a target GPU. * **Bug Fixes** * Improved Blackwell architecture handling, including support for devices with nonstandard architecture identifiers. * Prevented unsuitable CuTe DSL backends from being selected automatically when architecture support is unavailable. <!-- end of auto-generated comment: release notes by coderabbit.ai --> Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
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
7 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
Adds SM107 (compute capability 10.7) support to FlashInfer, so its existing kernels and backends compile and run on SM107 GPUs.
Key changes:
🔍 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
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
use_fp16_softmaxanduses_spcompresscontrols for paged and ragged attention variant selection.