fix(gemm): exclude cuDNN 9.23.0 SM90 split-k bf16 plans from the autotune space - #3974
YangXu1990uiuc wants to merge 1 commit into
Conversation
…tune space cuDNN 9.23.0 (backend_version 92300) miscomputes bf16 GEMM/BMM split-k plans on SM90/Hopper -- the same split-k output-layout bug that _cudnn_bf16_gemm_usable_or_skip already hard-bans for fp16 output. For bf16 output the DEFAULT plan is correct, so the envelope ban would be wrong-shaped: only the split-k plans (knob k17 = CUDNN_KNOB_TYPE_SPLIT_K_SLC) are broken, and they only execute when autotune(True) selects one -- which it does, because split-k wins the timing race on tall-K shapes and then silently returns garbage (found by the flashinfer-ai#3539 fuzzer's autotune-ON winner validation: mm_bf16 m63 n32 k2688, ratio 0.98 vs 0.0022 at the default plan). Per-plan enumeration on H100 + 9.23.0.39: exactly the five eng7_k17=4_* plans fail (ratio ~1.4), everything else is correct; the same matrix on 9.23.1.3 / 9.23.2.1 is all-correct, so the gate is exactly ==92300. 9.23.0.39 is a public PyPI nvidia-cudnn-cu12/cu13 release, so the window is reachable by users. Fix: CudnnBf16GemmRunner.get_valid_tactics drops plans whose structured name contains 'k17=' when backend_version()==92300 on SM90 (plan names need cudnn-frontend >= 1.25; without the API it falls back to the full list, where the fuzzer's numeric-only ledger entry still catches a bad winner). Default (non-autotuned) execution and every other arch/version are untouched. Verified on H100 + cuDNN 9.23.0.39: offered plans 15 -> 10 (all k17 plans dropped), 6/6 fresh autotune runs correct (worst ratio 0.0026, previously intermittently 1.39), fuzzer seed 1834712400 passes 3/3 (previously xfail); on SM100 the gate stays inactive (15/15 offered). AI-assisted (per-plan bisect + fix by Claude Code). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
|
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: defaults Review profile: CHILL Plan: Pro Run ID: 📒 Files selected for processing (1)
📝 WalkthroughWalkthroughThe cuDNN tactic search now filters ChangescuDNN tactic filtering
Estimated code review effort: 2 (Simple) | ~10 minutes Suggested reviewers: 🚥 Pre-merge checks | ✅ 5✅ Passed checks (5 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}
There was a problem hiding this comment.
Code Review
This pull request modifies get_valid_tactics in flashinfer/gemm/gemm_base.py to filter out broken split-k plans (k17=) for cuDNN version 9.23.0 on SM90 architectures to prevent silent corruption during autotuning. The reviewer suggested guarding the cudnn.backend_version() check with CUDNN_AVAILABLE to avoid potential NameError or linter issues when cuDNN is not installed.
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.
| if ( | ||
| cudnn.backend_version() == 92300 | ||
| and get_compute_capability(a.device)[0] == 9 | ||
| and hasattr(graph, "get_plan_name_at_index") | ||
| ): |
There was a problem hiding this comment.
To prevent potential NameError or static analysis/linter warnings when cudnn is not installed/available in the environment, we should guard the cudnn.backend_version() check with CUDNN_AVAILABLE, similar to how it is done in _cudnn_bf16_gemm_usable_or_skip.
if (
CUDNN_AVAILABLE
and cudnn.backend_version() == 92300
and get_compute_capability(a.device)[0] == 9
and hasattr(graph, "get_plan_name_at_index")
):
|
don't merge yet, will wait for 3707, then this need to be rewritten |
|
Closed as superseded by #3707: its |
2 participants
What
CudnnBf16GemmRunner.get_valid_tacticsnow excludes cuDNN split-k execution plans (structured plan name containingk17=, i.e.CUDNN_KNOB_TYPE_SPLIT_K_SLC) from the autotuning space when running on SM90 with cuDNN 9.23.0 exactly (backend_version() == 92300). One file, ~15 lines.Why
cuDNN 9.23.0 miscomputes bf16 GEMM/BMM split-k plans on SM90/Hopper — the same split-k output-layout bug that
_cudnn_bf16_gemm_usable_or_skip(landed in #3539) already hard-bans for fp16 output, where even the default plan is broken. For bf16 output the default plan is correct, so a wholesale ban would kill a working path; the broken plans only execute whenautotune(True)picks one — which it does, because split-k wins the timing race on tall-K shapes, and then silently returns garbage.Found by the #3539 fuzzer's autotune-ON winner validation (
mm_bf16m63 n32 k2688: autotuned ratio 0.98 vs 0.0022 at the default plan). Full analysis with the per-plan matrix: #3539 (comment).Per-plan enumeration on H100 + 9.23.0.39: exactly the five
eng7_k17=4_*plans fail (ratio ~1.4); every other plan is correct. The same matrix on 9.23.1.3 and 9.23.2.1 is all-correct, so the version gate is exactly==92300. Note 9.23.0.39 is a public PyPInvidia-cudnn-cu12/cu13release, so the silent-garbage window is reachable by real users.Why tactic-level (not an envelope ban)
backend="cudnn"users keep a working backend instead of a raise.autotune(True)still tunes — just among the correct plans.Plan names require cudnn-frontend ≥ 1.25 (
get_plan_name_at_index); without the API the filter falls back to the full list, where the #3539 fuzzer's numeric-only ledger entry still catches a bad winner. This is also a concrete argument for #3707's structured plan names:eng7_k17=4is a stable identifier where an integer plan index is version-fragile.Validation (H100 + cuDNN 9.23.0.39)
k17=plans dropped (names verified).autotune(True)runs correct (worst ratio 0.0026; previously intermittently 1.39).1834712400passes 3/3 deterministically (previously xfail-on-trip).🤖 Generated with Claude Code
Summary by CodeRabbit