Yanqinz/autotuner tactic - #3707
Merged
Merged
Conversation
yanqinz2
requested review from
aleozlx,
bkryu,
cyx-6,
dhiraj113,
jiahanc,
jimmyzho,
kahyunnam,
nv-yunzheq,
qsang-nv,
saltyminty,
samuellees,
sricketts,
yongwww,
yyihuang and
yzh119
as code owners
June 23, 2026 22:19
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:
📝 WalkthroughWalkthroughRefactors cuDNN GEMM execution to select plans through tactic-to-plan_index mapping, updates autotuner tactic handling and cache loading, raises the cuDNN frontend minimum version, and adjusts related tests and warning-based fallback paths. ChangescuDNN tactic→plan_index refactor
Estimated code review effort🎯 4 (Complex) | ⏱️ ~60 minutes Suggested labels
Suggested reviewers
Poem
🚥 Pre-merge checks | ✅ 3 | ❌ 2❌ Failed checks (1 warning, 1 inconclusive)
✅ Passed checks (3 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 |
![gemini-code-assist[bot]](https://avatars.githubusercontent.com/in/956858?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
Code Review
This pull request refactors the cuDNN GEMM backend to support fine-grained tuning tactics (engine/knob configurations) during graph compilation and execution, updates the minimum cuDNN backend version requirement for override-shape GEMM to 9.23.1, and bumps the nvidia-cudnn-frontend dependency to >=1.25.0. The review feedback suggests caching the resolved plan index in _get_cudnn_plan_index_for_tactic to eliminate runtime lookup overhead, and wrapping _is_cudnn_override_shape_available in a try...except block to prevent unhandled exceptions if cuDNN queries fail.
Important
The consumer version of Gemini Code Assist on GitHub is being sunset. Starting June 18, 2026, new organization installations will be blocked, and all code review activity will officially cease on July 17, 2026.
For more details on the timeline and next steps, please review the Help Documentation.
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
Actionable comments posted: 2
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@flashinfer/gemm/gemm_base.py`:
- Around line 2275-2280: The code currently allows any non-negative integer
tactic to be used as a concrete plan index, which permits legacy autotune caches
with unstable fixed-plan-index semantics to silently continue operating. Add
validation to reject non-negative integer tactics as legacy/invalid by checking
if tactic is an integer type and if so (and not equal to -1), set plan_index to
-1 instead of using the tactic value directly. This ensures old cached integer
plan indices are not silently accepted and forces the use of stable semantics.
- Around line 2696-2698: The override-shape FP4 graph builder is not applying
the same `eng0` engine deselection guard that exists in the fixed-shape builder.
Apply the same logic that excludes `eng0` when alpha or global scaling is
present and cuBLAS FP4-in-cuDNN is unavailable to the override-shape builder's
finalization step. Ensure that before calling `_finalize_cudnn_graph_for_tactic`
in the override-shape path, you check the same conditions as the fixed-shape
builder and apply the necessary engine deselection to prevent autotuning from
selecting unsupported engines.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: defaults
Review profile: CHILL
Plan: Pro
Run ID: c083a14e-d6f9-4f16-8c0c-db0bc9c10a1c
📒 Files selected for processing (3)
flashinfer/gemm/gemm_base.pyflashinfer/gemm/gemm_bf16_fp4_cudnn.pyrequirements.txt
Contributor
There was a problem hiding this comment.
🧹 Nitpick comments (1)
flashinfer/autotuner.py (1)
425-425: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winAlign remaining tactic contracts with the new
AnymodelAfter widening
forward/choose_onetoAny, there are still integer-specific remnants (TunableRunner.forwarddoc text at Line 433 andsearch_cachereturn annotation at Line 947) that implytacticis alwaysint. Please update those to an opaque tactic type consistently to avoid type-checking drift and misleading API docs.Also applies to: 1102-1102, 1113-1115
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@flashinfer/autotuner.py` at line 425, Update the tactic-related documentation and type annotations throughout the file to consistently use the `Any` type model. Specifically, update the docstring for the TunableRunner.forward method to remove language implying tactic is always an integer, change the return type annotation of the search_cache method from int to Any to match the new tactic contract, and similarly update any other type annotations or documentation at the additional locations (around lines 1102 and 1113-1115) that still reference tactic as an integer type to instead reflect the opaque `Any` tactic model.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Nitpick comments:
In `@flashinfer/autotuner.py`:
- Line 425: Update the tactic-related documentation and type annotations
throughout the file to consistently use the `Any` type model. Specifically,
update the docstring for the TunableRunner.forward method to remove language
implying tactic is always an integer, change the return type annotation of the
search_cache method from int to Any to match the new tactic contract, and
similarly update any other type annotations or documentation at the additional
locations (around lines 1102 and 1113-1115) that still reference tactic as an
integer type to instead reflect the opaque `Any` tactic model.
ℹ️ Review info
⚙️ Run configuration
Configuration used: defaults
Review profile: CHILL
Plan: Pro
Run ID: 6788c460-c4dc-4acf-9309-9e679bd5390a
📒 Files selected for processing (2)
flashinfer/autotuner.pyflashinfer/gemm/gemm_base.py
🚧 Files skipped from review as they are similar to previous changes (1)
- flashinfer/gemm/gemm_base.py
yanqinz2
force-pushed
the
yanqinz/autotuner-tactic
branch
from
June 25, 2026 18:53
f359f89 to
8a37e5b
Compare
Contributor
There was a problem hiding this comment.
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (2)
flashinfer/gemm/gemm_bf16_fp4_cudnn.py (1)
510-527: 🩺 Stability & Availability | 🟠 Major | ⚡ Quick winDo not retry default fixed-shape failures as fallback.
If
_use_override_shapeis false andtactic == -1, this catch block retries the same fixed-shape default path after warning, which can duplicate OOM/validation failures and obscure the real error. Re-raise in that case; keep the fallback for non-default tactics or override-shape recovery.Suggested guard
except Exception as exc: + if tactic == -1 and not self._use_override_shape: + raise warnings.warn( "cuDNN bf16-fp4 GEMM tactic failed; falling back to default " f"tactic=-1. ({exc})",🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@flashinfer/gemm/gemm_bf16_fp4_cudnn.py` around lines 510 - 527, The fallback in the cuDNN bf16-fp4 GEMM path is retrying the same fixed-shape default tactic when `_use_override_shape` is false and `tactic == -1`, which can repeat the original failure instead of surfacing it. Update the `except Exception as exc` block in `build_cudnn_bf16_fp4_graph` / the surrounding tactic-selection logic to re-raise immediately for that default fixed-shape case, and only keep the warning-plus-retry behavior for non-default tactics or when recovering via `_use_override_shape`.Source: Linters/SAST tools
flashinfer/gemm/gemm_base.py (1)
3478-3493: 🩺 Stability & Availability | 🟠 Major | ⚡ Quick winAvoid retrying the same default fixed-shape failure.
When
_use_override_shapeis false andtactic == -1, these handlers catch a default execution failure, warn that they are “falling back” to the same default path, then retry it. Re-raise in that case, and reserve fallback for cached/non-default tactics or override-shape-to-fixed-shape recovery.Suggested guard
except Exception as exc: + if tactic == -1 and not self._use_override_shape: + raise warnings.warn( "... falling back to default tactic=-1. ...", stacklevel=2, )Also applies to: 3992-3998, 4970-4984, 5441-5458, 8788-8802
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@flashinfer/gemm/gemm_base.py` around lines 3478 - 3493, In the GEMM fallback handlers in _cudnn_gemm_fp8 and the related retry paths, avoid re-invoking the same default fixed-shape execution when _use_override_shape is false and tactic == -1. Add a guard before the fallback call so these blocks re-raise the original exception in that case, and only retry for non-default/cached tactics or when recovering from override-shape to fixed-shape. Apply the same logic consistently across the matching exception handlers in the other GEMM entry points.Source: Linters/SAST tools
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Outside diff comments:
In `@flashinfer/gemm/gemm_base.py`:
- Around line 3478-3493: In the GEMM fallback handlers in _cudnn_gemm_fp8 and
the related retry paths, avoid re-invoking the same default fixed-shape
execution when _use_override_shape is false and tactic == -1. Add a guard before
the fallback call so these blocks re-raise the original exception in that case,
and only retry for non-default/cached tactics or when recovering from
override-shape to fixed-shape. Apply the same logic consistently across the
matching exception handlers in the other GEMM entry points.
In `@flashinfer/gemm/gemm_bf16_fp4_cudnn.py`:
- Around line 510-527: The fallback in the cuDNN bf16-fp4 GEMM path is retrying
the same fixed-shape default tactic when `_use_override_shape` is false and
`tactic == -1`, which can repeat the original failure instead of surfacing it.
Update the `except Exception as exc` block in `build_cudnn_bf16_fp4_graph` / the
surrounding tactic-selection logic to re-raise immediately for that default
fixed-shape case, and only keep the warning-plus-retry behavior for non-default
tactics or when recovering via `_use_override_shape`.
ℹ️ Review info
⚙️ Run configuration
Configuration used: defaults
Review profile: CHILL
Plan: Pro
Run ID: ba169850-bd5f-431d-925f-a2f2a1635afc
📒 Files selected for processing (5)
flashinfer/autotuner.pyflashinfer/gemm/gemm_base.pyflashinfer/gemm/gemm_bf16_fp4_cudnn.pyrequirements.txttests/autotuner/test_autotuner_bmm_fp8.py
🚧 Files skipped from review as they are similar to previous changes (2)
- requirements.txt
- flashinfer/autotuner.py
Collaborator
|
Overall, the engine/knob tactic representation and the cuDNN version gating look reasonable. I have the following requested changes/hardening suggestions: 1. Fix the MXFP8 test so it actually exercises the override-shape path
mat2 = torch.randn([b, n, k]).transpose(-2, -1).contiguous()
mat2_q, mat2_scale = mxfp8_quantize(mat2, ...)This produces a row-major logical It should instead be constructed as: weight = torch.randn([b, n, k], device="cuda", dtype=input_dtype)
weight_q, weight_scale = mxfp8_quantize(weight, is_sf_swizzled_layout)
B = weight_q.transpose(-2, -1)
reference = torch.bmm(input_mat, weight.transpose(-2, -1))Please also assert the expected stride, and ideally ensure that no fallback warning is emitted. The current test passes numerically only because the runner silently falls back to the fixed-shape path, so it does not cover the new override-shape execution path. This test bug predates this PR, but the new fallback behavior makes it especially easy for the test to pass without exercising the intended code. 2. Do not treat a failed tactic's fallback execution as a successful tactic during profilingFallback is desirable for inference robustness, but during autotuning, a failed tactic should be reported back to the autotuner as failed rather than being timed as Otherwise, the autotuner may persist a tactic that never executed and actually measured the fallback implementation. A preferable policy is:
This preserves serving reliability while ensuring that tuning results represent the requested tactics. 3. Include BF16 matrix layout in the persistent cache keyThe in-memory runner identity distinguishes the A/B layout flags, but the persisted file-cache key currently does not distinguish: is_a_k_major
is_b_k_majorAs a result, different NT/TN/layout configurations can reuse a tactic tuned for another layout after serialization. Please include these flags in This may be uncommon, but cache-key bugs are extremely difficult to diagnose because they produce stable, configuration-dependent performance regressions. 4. Version the persisted cuDNN tactic representationThis PR changes cuDNN tactics from integer plan indices to engine/knob descriptions. Please add a tactic-schema/cache-format version and invalidate or explicitly reject legacy non-negative integer cuDNN tactics. FlashInfer version metadata helps for normal releases, but it does not fully protect development builds, same-version branch changes, manually copied configs, or wildcard metadata. A persisted plan index must never be interpreted as a tactic under the new representation. |
yanqinz2
force-pushed
the
yanqinz/autotuner-tactic
branch
from
June 30, 2026 00:53
8a37e5b to
1e781d0
Compare
Contributor
There was a problem hiding this comment.
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (2)
flashinfer/gemm/gemm_base.py (1)
3478-3493: 🎯 Functional Correctness | 🟠 Major | 🏗️ Heavy liftDon’t fallback during autotuner profiling.
These
except Exceptionblocks make a failed candidate tactic return successfully after executingtactic=-1, so autotuning can cache an unsupported tactic as successful. Gate this fallback to normal inference only; during profiling/tuning, re-raise so the autotuner marks the candidate failed and only uses fallback after all candidates fail.Also applies to: 3997-4003, 4975-4989, 5446-5464, 8798-8813
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@flashinfer/gemm/gemm_base.py` around lines 3478 - 3493, The cuDNN fp8 GEMM exception handlers currently swallow candidate failures by retrying with tactic=-1, which can make autotuner profiling treat unsupported tactics as successful. Update the relevant fallback paths in the affected GEMM helpers (including the code around _cudnn_gemm_fp8 and the other listed except Exception blocks) to distinguish normal inference from autotuner profiling/tuning, and only apply the tactic=-1 fallback during normal inference. During profiling, re-raise the exception so the autotuner records the candidate as failed and only relies on fallback after all candidates fail.Source: Linters/SAST tools
flashinfer/gemm/gemm_bf16_fp4_cudnn.py (1)
510-538: 🎯 Functional Correctness | 🟠 Major | 🏗️ Heavy liftDon’t fallback during autotuner profiling.
This catches candidate tactic failures and then succeeds via fixed-shape
tactic=-1, which can cause autotuning to record the failed candidate as supported. Re-raise in tuning/profiling mode and keep the warning fallback only for normal inference.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@flashinfer/gemm/gemm_bf16_fp4_cudnn.py` around lines 510 - 538, The fallback inside the cuDNN bf16-fp4 GEMM path should not run during autotuner profiling, because the fixed-shape tactic=-1 execution can make a failed candidate look supported. Update the exception handling around the cuDNN graph execution in the bf16_fp4 GEMM flow to re-raise when tuning/profiling is active, and keep the warning plus tactic=-1 fallback only for normal inference. Use the existing cuDNN GEMM helpers and control flow in gemm_bf16_fp4_cudnn.py, especially the block that calls build_cudnn_bf16_fp4_graph and execute_cudnn_bf16_fp4_graph.Source: Linters/SAST tools
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Outside diff comments:
In `@flashinfer/gemm/gemm_base.py`:
- Around line 3478-3493: The cuDNN fp8 GEMM exception handlers currently swallow
candidate failures by retrying with tactic=-1, which can make autotuner
profiling treat unsupported tactics as successful. Update the relevant fallback
paths in the affected GEMM helpers (including the code around _cudnn_gemm_fp8
and the other listed except Exception blocks) to distinguish normal inference
from autotuner profiling/tuning, and only apply the tactic=-1 fallback during
normal inference. During profiling, re-raise the exception so the autotuner
records the candidate as failed and only relies on fallback after all candidates
fail.
In `@flashinfer/gemm/gemm_bf16_fp4_cudnn.py`:
- Around line 510-538: The fallback inside the cuDNN bf16-fp4 GEMM path should
not run during autotuner profiling, because the fixed-shape tactic=-1 execution
can make a failed candidate look supported. Update the exception handling around
the cuDNN graph execution in the bf16_fp4 GEMM flow to re-raise when
tuning/profiling is active, and keep the warning plus tactic=-1 fallback only
for normal inference. Use the existing cuDNN GEMM helpers and control flow in
gemm_bf16_fp4_cudnn.py, especially the block that calls
build_cudnn_bf16_fp4_graph and execute_cudnn_bf16_fp4_graph.
ℹ️ Review info
⚙️ Run configuration
Configuration used: defaults
Review profile: CHILL
Plan: Pro
Run ID: c4dc52c9-7e34-4cd5-ae24-fb25ded3fa81
📒 Files selected for processing (6)
flashinfer/autotuner.pyflashinfer/gemm/gemm_base.pyflashinfer/gemm/gemm_bf16_fp4_cudnn.pyrequirements.txttests/autotuner/test_autotuner_bmm_fp8.pytests/gemm/test_bmm_mxfp8.py
🚧 Files skipped from review as they are similar to previous changes (2)
- requirements.txt
- flashinfer/autotuner.py
Collaborator
Author
|
Regarding above comment from @YangXu1990uiuc
|
Contributor
There was a problem hiding this comment.
🧹 Nitpick comments (1)
flashinfer/autotuner.py (1)
385-385: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winFinish widening the tactic contract beyond
List/integer wording.
get_valid_tactics()now supports arbitrary tactic values, butList[Any]is still too concrete for tuple-returning runners, andforward()still documentstacticas an integer ID. PreferSequence[Any]and runner-defined wording so cuDNN engine/knob tuple tactics match the public contract.Suggested contract cleanup
- ) -> List[Any]: + ) -> Sequence[Any]: @@ - tactic: Integer ID specifying which implementation tactic to use. + tactic: Runner-defined selector specifying which implementation tactic to use.If
Sequenceis not already imported:-from typing import Any, ... +from typing import Any, Sequence, ...Also applies to: 425-435
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@flashinfer/autotuner.py` at line 385, The tactic contract is still too narrow in both the return type and the `forward()` docs: update `get_valid_tactics()` to use `Sequence[Any]` instead of `List[Any]`, and revise `forward()`’s `tactic` wording so it describes runner-defined tactic values rather than an integer ID. Make sure the `autotuner` public API consistently allows tuple-style cuDNN engine/knob tactics, and import `Sequence` if needed.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Nitpick comments:
In `@flashinfer/autotuner.py`:
- Line 385: The tactic contract is still too narrow in both the return type and
the `forward()` docs: update `get_valid_tactics()` to use `Sequence[Any]`
instead of `List[Any]`, and revise `forward()`’s `tactic` wording so it
describes runner-defined tactic values rather than an integer ID. Make sure the
`autotuner` public API consistently allows tuple-style cuDNN engine/knob
tactics, and import `Sequence` if needed.
yanqinz2
force-pushed
the
yanqinz/autotuner-tactic
branch
from
July 15, 2026 16:55
1990cae to
382c1b6
Compare
YangXu1990uiuc
added a commit
that referenced
this pull request
Jul 16, 2026
Quantize the contiguous [b,n,k] weight and pass the TRANSPOSE VIEW as B (column-major [b,k,n], K contiguous), instead of .contiguous()-ing the transpose before quantization. The old layout is rejected by the cuDNN override-shape path's stride check; CI never saw it only because its older cudnn-frontend (<1.24) keeps the override path disabled -- the moment the frontend floor rises (e.g. #3707 pins >=1.25.0) every cudnn case here would fail. Matches the unified fuzzer's bmm_mxfp8 adapter recipe. Verified on SM100 + cuDNN 9.24.0.43 (torch main's current pin): 3 cudnn cases pass (previously ValueError), 2 cutlass cases skip (SM12x-only). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
5 tasks
Collaborator
|
Following up on review point #2 from this PR ("do not treat a failed tactic's fallback execution as a successful tactic during profiling"), which was deferred here as potentially needing a broad autotuner change: it turned out to be a local change — |
Saddss
added a commit
to Saddss/flashinfer
that referenced
this pull request
Jul 27, 2026
…navailable On SM12x (RTX 5090) with override_shape unavailable (cuDNN backend < 9.23.1), cuDNN's per-shape policy=ALL execution-graph build is unbounded host-side work at serving time (RFC flashinfer-ai#3920, runner-contract rule 6), and the cuDNN FP8 serving path raises an async cudaErrorMisalignedAddress that flashinfer-ai#3707's synchronous forward() try/except cannot catch. Gate cuDNN out of the SM12x auto candidate list in that case; SM100/103/107/110 and SM12x-with-override are unchanged (cuDNN stays a candidate where M-bucketing amortizes the build). Verified on RTX 5090 (SM120, vLLM nightly, online-replay load test), same container, only the FlashInfer wheel differs: - backend 92000 (override unavailable) + cuDNN in auto: ~5 rounds crash (misaligned), p50 ~17s - backend 92000 + this gate (cuDNN removed): 12/12 PASS, p50 10.31s, 0 crashes - backend 92400 (override available, nvidia-cudnn-cu13 9.24) + cuDNN in auto, NO gate: 12/12 PASS, p50 10.69s, 0 crashes -- cuDNN is safe when override_shape is available, validating the conditional gate's else branch. The gate removes only cuDNN; cutlass_sm12x/cublas kernels and within-backend winner selection are unchanged, so eliminating the 92000 crash by removing cuDNN identifies cuDNN as the faulting path. (An earlier nvidia-cudnn-cu12 9.24 attempt was a CUDA 12/13 library mismatch, not a cuDNN signal.) Signed-off-by: Saddss <2872669061@qq.com>
5 tasks
bkryu
added a commit
that referenced
this pull request
Aug 6, 2026
…navailable (#4165) ## 📌 Description Gate the cuDNN FP8 GEMM backend out of the `bmm_fp8(backend="auto")` candidate list on SM12x when cuDNN `override_shape` is unavailable (cuDNN backend < 9.23.1). On this narrow configuration the cuDNN path is both unsafe and unnecessarily slow; cublas/cutlass already cover the workload. ## Root cause Two independent failure modes on SM12x (RTX 5090), both removed by the gate: 1. **Unbounded host-side compilation (perf).** Without `override_shape`, cuDNN builds a fresh `policy=ALL` execution graph per distinct serving `(M,K,N)` shape (~250 ms host-side; measured 236 ms/shape). Under load this is unbounded — RFC #3920, runner-contract rule 6. 2. **Async CUDA fault (crash).** The cuDNN FP8 serving path raises `cudaErrorMisalignedAddress` asynchronously at the next sync, which #3707's `forward()` try/except cannot catch (it only handles synchronous exceptions). ## 📊 Evidence chain (RTX 5090, SM120, vLLM nightly, online-replay load test) Same container; only the FlashInfer wheel (and cuDNN package) differs. **Reproducibility (input/output lengths only)** - Model checkpoint (open-sourced): <https://huggingface.co/wangqia0309/gemma-4-26B-A4B-it_nvfp4_experts_fp8_dense_fp8_attn-kv_fp8> (NVFP4 mixed-precision: FP8 dense + attention, FP8 KV cache) - Input: `max-model-len 8192`, `max-num-batched-tokens 8192`, `max-num-seqs 64` - Output: `max-tokens 200` - MTP: `num_speculative_tokens=3` - vLLM dtype/quant config: `--quantization modelopt --kv-cache-dtype fp8 --moe-backend cutlass -O3 --gpu-memory-utilization 0.95` | FI build | cuDNN in SM12x auto | override_shape | result | |---|---|---|---| | 0.6.15.post1 (BAD baseline) | yes | unavailable (backend 92000) | 6 rounds then crash; p50 96/78/71/78/74/40 s | | main + #3707 | yes | unavailable (backend 92000) | ~5 rounds crash (`misaligned`); p50 26/17.6/17/17.5 s | | **main + #3707 + this gate** | no (gated) | unavailable (backend 92000) | **12/12 PASS, steady p50 10.31 s, 0 crashes** | | main + #3707 (NO gate) | yes | **available** (backend 92400, `nvidia-cudnn-cu13` 9.24) | **12/12 PASS, steady p50 10.69 s, 0 crashes** | The last row is the **clean** override-available test: with the cu13 cuDNN 9.24 stack (override_shape available), cuDNN-in-auto does *not* crash and matches the gated run's p50 — validating the conditional gate's `else` branch (keep cuDNN where override_shape amortizes the build). **Why the gate proves the faulting path is cuDNN (elimination).** The gate removes *only* cuDNN from the candidate list; the `cutlass_sm12x` and `cublas` kernels and the autotuner's within-backend winner selection are unchanged. At serving time only the winning runner executes, and on 92000-without-override the cuDNN runner is selected (per-tactic warm median ≈ 0.05–0.07 ms, winning 20/20 in the §7.3 micro-bench). Removing cuDNN eliminates the crash; the remaining cutlass/cublas kernels do not fault (12/12). Therefore the faulting path is cuDNN — no `compute-sanitizer` pinning needed. **Independent reproduction:** handoff §6 — on post2, removing cuDNN from the SM120 candidate list → 12/12, steady p50 ≈ 12 s. ## Scope and limitation - **Scoped fix.** The gate is `is_sm120_supported and not _is_cudnn_override_shape_available()` — SM100/103/107/110 and SM12x-with- override are unchanged (cuDNN stays a candidate where M-bucketing amortizes the build). - **"Upgrade cuDNN" is tested clean here.** With the matching CUDA-13 cuDNN (`nvidia-cudnn-cu13==9.24.0.43`, backend 92400, override_shape available) + cuDNN still in `auto`, the run is **12/12 PASS, p50 10.69 s, 0 crashes** — cuDNN is safe when `override_shape` is available, so the conditional gate keeps it. (An earlier `nvidia-cudnn-cu12==9.24` attempt was a CUDA 12/13 library mismatch that broke cuBLASLt symbol loading — not a cuDNN-FP8 signal.) - **No upstream improvement yet:** vLLM v0.26.0 stable pins `flashinfer-python==0.6.14` (BAD per the §4 bisect); vLLM nightly pins 0.6.15.post1 (BAD); FlashInfer main + #3707 still crashes on SM12x without this gate. ## 🔍 Related Issues - RFC #3920 (Autotuner v2, runner-contract rule 6) — unbounded serving-time compilation - #3707 (engine/knob tactics + sync fallback) — fixes stale plan-index crash, not the async fault - #3437 (multi-tactic / `policy=ALL` for cuDNN) — introduced the per-shape ALL build cost - #2914 (cuBLASLt + FP8 multi-tactic autotuning) — first BAD boundary - #3566, #3673, #3255 — SM12x cuDNN/plan-index field reports ## 🚀 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.). Added `test_bmm_fp8_heuristic_gates_cudnn_on_sm12x_without_override_shape` — monkeypatches SM version + cuDNN availability so it runs on any host without SM12x hardware or a cuDNN install; asserts cuDNN is excluded when `override_shape` is unavailable and retained when available. ## Reviewer Notes - Scope is deliberately minimal: one heuristic branch + one test. No change to cuDNN code, the override_shape path, or non-SM12x architectures. - The one-shot warning uses a module-level `set` (mutation, no `global`) so it does not spam per-call (the heuristic runs on every `bmm_fp8(backend="auto")`). - DCO: `Signed-off-by: Saddss <2872669061@qq.com>`. Signed-off-by: Saddss <2872669061@qq.com> Co-authored-by: Brian K. Ryu <bryu@nvidia.com>
aleozlx
added a commit
that referenced
this pull request
Sep 5, 2026
…ployment-matched measurement, runner contract (#3861) ## 📌 Description **Autotuner v2** — an evolution of the autotuner exposed through a standalone `autotune_v2()`, additive over `autotune()` (which stays byte-identical). It started as a managed persistent cache and has grown, across review, into five composing pillars: 1. **Managed persistence** — FlashInfer owns cache keys, schema, environment isolation, and invalidation; frameworks stop hashing filenames, parsing the JSON schema, broadcasting raw cache files, and managing staleness. Per-entry atomic store, safe for concurrent all-ranks tuning. 2. **Deployment-matched measurement** (`MeasurementPolicy`) — tune the way you serve (eager vs CUDA-graph host-cost semantics); part of the store's environment identity. 3. **Runner contract** — self-describing tactics + runtime revalidation (`validate_tactic`, covering in-memory and on-disk hits) so a stale/invalid tactic is a loud fallback, never a silent wrong kernel or a dead server. 4. **Distributed consistency** — a shared store plus `autotune_v2_reload()` finalize step so homogeneous ranks converge on byte-identical tactics (composes with #3187). 5. **Selection quality** — a default-candidate coverage guard (a tuned result cannot lose to not-tuning at the probe inputs) and an accuracy harness that quantifies regret against a deployment-faithful oracle. > **RFC** (design + evidence, for public/customer review): #3920. Detailed rationale, the review-round fixes, framework-migration drafts, and the production-B200 validation are in the thread below. ## Summary FlashInfer owns the persistence and compatibility rules for autotuning results, so frameworks stop constructing cache filenames, parsing FlashInfer's JSON schema, broadcasting raw cache files, and managing invalidation themselves — and, via `MeasurementPolicy`, tuning measures the way the deployment actually runs. Exposed as a **standalone `autotune_v2()` context manager**, deliberately disjoint from `autotune(cache=<json path>)` (v1) so the design can iterate without inheriting v1 semantics — `flashinfer/autotuner/` changes are purely additive hooks and `autotune()` itself is untouched. ## API ```python with flashinfer.autotune_v2(measure=MeasurementPolicy(execution_mode="cuda_graph")): model(dummy_inputs) # startup warmup: tune + publish atomically model(inputs) # serving: reuses entries, no context needed # fresh process, same environment: hydrate + serve with flashinfer.autotune_v2(mode="replay", measure=policy): pass ``` | `mode` | `persistent_cache` | meaning | |---|---|---| | `"tune"` | True *(default)* | tune misses, publish to disk | | `"tune"` | False | tune in-memory only (no disk I/O) | | `"replay"` | True | serve from the on-disk cache (no profiling) | | `"replay"` | False | no-op unless already hydrated | **Attach semantics**: `persistent_cache=True` attaches the store for the remainder of the process (like v1's `load_configs`) — a survey of vLLM/sglang showed both serve **outside** any context, so a context-scoped store would silently regress serving to heuristics after warmup exits. The context scopes only *when profiling cost may be paid*. `cache_root` (or `FLASHINFER_AUTOTUNE_CACHE_DIR`) is a placement-only **directory**; the schema/environment namespaces live below it. ## MeasurementPolicy — tune the way you deploy The decisive question is whether per-call **host** cost counts: under CUDA-graph serving it is paid once at capture (excluding it is correct); under eager serving it is paid every call. Measured on SM100 (`bmm_fp8` M=8, kernel ~8µs): the same cuDNN candidate reads **8.0µs host-excluded vs ~330µs host-included** — the eager ranking flips to cublas < cutlass ≪ cuDNN, matching real framework behavior. ```python MeasurementPolicy(execution_mode="cuda_graph") # profile under real capture+replay MeasurementPolicy(execution_mode="eager") # host cost included (no delay kernel) # default "auto" = today's legacy behavior; cold_l2 composes. The timer is derived # from execution_mode and is always CUDA events: "auto"/"cuda_graph" use the # events+delay-kernel window, "eager" uses events without the delay kernel. # CUPTI activity-span timing is NOT selected by any public mode; it exists only # as a private diagnostic override (_timer="cupti") used by the accuracy harness. ``` The policy is part of the store's **environment identity** (manifest → env hash): entries tuned under different policies never overwrite each other. A CUPTI activity-span timer is implemented as a **diagnostic-only** route (per-iteration GPU spans, median; no delay kernel; `cuptiIsTracingSessionRunning` pre-check + graceful events fallback around the legacy single-subscriber limitation) and is reachable only via the private `_timer="cupti"` override, not through `execution_mode`. It is not used for any shipped policy; see the validation report's CUPTI appendix for why (it cannot rank eagerly-launched multi-kernel ops such as `moe_cute_dsl`). ## Store design ``` <root>/v2/<environment_hash>/ # sha256 of canonical manifest (flashinfer/CUDA/ manifest.json # cuBLAS/cuDNN backend+frontend/GPU [+ policy]) entries/<operation_hash>.json # one atomic file per tuned operation ``` - One entry per tuned op, published via tempfile + atomic `os.replace` as soon as it is tuned — no exit-time read/merge/write, no locks; concurrent all-ranks tuning is safe (redundant work, last valid write wins), so per-rank cache files become unnecessary. - Missing/malformed/key-mismatched entries are cache misses, never errors; each entry embeds its canonical key. - Hit/miss memoized (decoded once); the serving hot path re-reads neither the filesystem nor JSON. - Incompatible environments hash to different directories: nothing to invalidate manually, downgrade-safe. - v1 and v2 never share in-memory state — a v1 `save_configs` structurally cannot observe v2 entries (regression-tested). ## Also included - **Runner contract**: `validate_tactic(inputs, tactic)` revalidation on in-memory and on-disk hits (a rejected tactic is a loud fallback, not a blind replay); default-candidate coverage guard. - **Distributed**: `autotune_v2_reload()` finalize step for rank convergence. - `benchmarks/bench_autotuner_accuracy.py`: accuracy harness (oracle sweep → regret / top-1 / winner-flip-rate under each policy, scored against both oracles), with an interleaved drift-protected oracle and per-round SM-clock recording. **Validated on a production B200** across bmm_fp8 / mm_fp4 / cutlass MoE (thread). - Observability: loud INFO at store attach, once-per-op hit provenance logs. ## Explicitly out of scope (this PR) - opaque export/install artifact for multi-node distribution (only needed for leader-tunes-then-broadcast without a shared FS; under discussion) - flipping `execution_mode="auto"` to `"cuda_graph"` (pending forced-capture validation across the op suite) - per-runner adoption of `validate_tactic` (cuDNN first, via its #3707 structured tactics — the hook is here; runners opt in separately) - representative-probe construction for MoE (EP-deflated shapes / skewed-routing dummies — the #3622/#3537 op-level track; sibling PRs) - serialized backend artifacts (e.g. cuDNN plan sidecar, after #3707); GC/pruning; cross-job single-flight; comm-op tuning; any change to v1 behavior ## Relation to existing work Orthogonal to #3707 (structured `(engine_id, knobs)` cuDNN tactics): entries store whatever tactic representation runners produce — compound-tuple round-trip is covered by tests; the environment hash contains index-tactic drift meanwhile. ## Testing 238 GPU-free tests pass in `tests/autotuner/` (the v2 file `test_autotune_cache_v2.py` covering: both real consumer patterns — serve-after-context-exit, hydrate-only; corruption tolerance; env-hash and measurement-policy isolation; bare-serving-uses-ambient-identity and cache_root-switch repopulation; atomicity; v1 coexistence; runner-reorder / extras-collision key completeness; guard win-persist-replay; reload convergence; same-process revalidation fallback — plus the upstream suites incl. the #3187 AST guards). Pre-commit clean. On-GPU validation on SM100 and a **production B200**: tune → atomic publish → fresh-process serve; a 4-process concurrent tuning storm into one shared store (SGLang all-ranks pattern) with store integrity intact; and the accuracy harness across three ops (bmm_fp8 / mm_fp4 / cutlass MoE) — details, matrices, and framework-migration drafts in the thread. Branch went through several external review rounds plus a four-lens cleanup pass. 🤖 Generated with [Claude Code](https://claude.com/claude-code) ## 🔍 Related Issues RFC #3920 (design) · #3707 (structured tactics, composed) · #3187 (cross-rank sync, composed) · addresses classes from #3186, #3537, #3566, #3622, #3648, #3719 (see the RFC's issue-bucket table) ## 🚀 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. ## 🧪 Tests - [x] Tests have been added or updated as needed (`tests/autotuner/test_autotune_cache_v2.py`, 28 GPU-free tests; full `tests/autotuner/` suite green). - [x] All tests are passing. ## Reviewer Notes Opt-in and additive: `autotune()` and the v1 cache path are byte-identical; v2 state lives behind `autotune_v2()`. Suggested review order: `flashinfer/autotune_cache.py` (store + API), then the four hook sites in `flashinfer/autotuner/autotuner.py` (search_cache 2.5, publish, measurement routing, store stack), then the tests. <!-- note to self: claude::28440023-66ee-410c-a0aa-234997e886d7 — "Autotuner v2 report review" cwd /home/scratch.yanxu_libs/flashinfer · workspace /tmp/claude-25653/autotune-v2-rebase --> --------- Co-authored-by: Claude Fable 5 <noreply@anthropic.com> Co-authored-by: Alex Yang <aleyang@nvidia.com> Co-authored-by: Vincent Tombari <34876120+Vinnie6167@users.noreply.github.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
4 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.
This PR contains a series of changes to make cudnn GEMM backend stable during autotuning.
🚀 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
nvidia-cudnn-frontendto ≥ 1.25.0.