Skip to content

test(vllm-tensorizer): Bump flashinfer 0.6.14 -> 0.6.15 (candidate GB300 MoE hang fix)#188

Draft
nehashah-eng wants to merge 1 commit into
mainfrom
neha/flashinfer-0.6.15-gb300-moe-hang-test
Draft

test(vllm-tensorizer): Bump flashinfer 0.6.14 -> 0.6.15 (candidate GB300 MoE hang fix)#188
nehashah-eng wants to merge 1 commit into
mainfrom
neha/flashinfer-0.6.15-gb300-moe-hang-test

Conversation

@nehashah-eng

Copy link
Copy Markdown

Summary

Testing whether flashinfer 0.6.15 fixes an intermittent production hang, not a verified fix yet — draft on purpose.

  • We've hit a reproducible hang on dynamo-vllm (Kimi-K2.6, dflash speculative decoding + flashinfer_trtllm MoE, TP=4, GB300) across multiple engine versions, confirmed 3 times now with live py-spy/gdb captures. Every capture shows the same signature: EngineCore blocked waiting on an RPC response while 3-4 TP worker processes are simultaneously stuck mid-cuLaunchKernelEx/cuLaunchKernel inside flashinfer::trtllm_fp4_block_scale_moe's routing/quantization kernels, frozen (zero progress) for the full ~9-12 minute hang window before vLLM's own RPC timeout kills the engine.
  • flashinfer-ai/flashinfer#3971 documents a similar-sounding GB300 (SM103) hang in the TRTLLM_GEN_BMM kernel family, regressed between flashinfer 0.6.13 -> 0.6.14 (the version currently pinned in this repo). A flashinfer maintainer suggested PR #3973 (merged into v0.6.15) as a possible fix — but this is not confirmed even by the flashinfer team, and the exact stuck kernel/mechanism in #3971 (GPU-resident kernel stuck in an mbarrier wait, per their own core dumps) doesn't exactly match what we captured (CPU thread stuck inside the launch call itself). Treat this as a plausible candidate, not a known fix.

What this PR does

Bumps flashinfer-commit from v0.6.14 to v0.6.15 for both CUDA variants of the vllm-commit: '752a3a504485790a2e8491cacbb35c137339ad34' (v0.25.1) matrix entry in .github/configurations/vllm-tensorizer.yml — this is the exact vLLM ref currently deployed for the affected dynamo-vllm image.

Patch rebase (important — please review)

This repo carries a custom patch (flashinfer-v0.6.14-deepseekv3-no-group-routing.patch, originally by @aeldeib) fixing a GLM-5.2 NaN-collapse bug in the DeepSeekV3 no-group routing path. I rebased it onto v0.6.15 (renamed to flashinfer-v0.6.15-deepseekv3-no-group-routing.patch) and verified the result against the real v0.6.15 source (not just assumed):

  • Second hunk dropped (csrc/trtllm_fused_moe_runner.cu, routingData.mSumEpsilon = 1e-20f;): this no longer applies — upstream's own v0.6.14→v0.6.15 diff independently adds the identical line at the same call site with an equivalent explanatory comment. Confirmed by diffing the real v0.6.15 source directly; this is provably redundant now, not just a textual conflict.
  • First hunk kept (RoutingCustomPolicy.cuh, warp-per-token applyWithScores): still applies cleanly (git apply, verified against real v0.6.15 source, 1-line offset). Not verified: whether it's still functionally necessary/correct, since upstream also generalized the surrounding postprocess-dispatch architecture in this version (new kNeedsAux/applyWithAux/PolicyPairNeedsAux traits replacing the narrower special-casing this hunk patches around). @aeldeib — would appreciate your eyes on whether this hunk still does the right thing given that refactor, before this merges. Full reasoning is in the patch file's header comment.

Test plan

  • CI build succeeds (config-file change should trigger a new vllm-tensorizer build matrix per .github/workflows/vllm-tensorizer.yml's paths: trigger)
  • @aeldeib (or GLM-5.2 owner) confirms the patch rebase is correct, or provides an updated patch
  • Once a dynamo-vllm image is built from this base, re-run the same load-test harness that reproduced the hang (multiple sweeps — this bug is intermittent, 1 clean run doesn't confirm a fix) and confirm the hang no longer occurs
  • If it does still occur, capture is still armed/documented (~/Projects/cognition/results-v025/) to compare the stack signature against pre-bump captures

Not planning to merge this until the above is resolved — draft is intentional.

🤖 Generated with Claude Code

…E hang fix)

Testing candidate: flashinfer-ai/flashinfer#3971 documents an intermittent
GB300 (SM103) hang in the TRTLLM_GEN_BMM-family kernel used by
flashinfer_trtllm MoE, regressed between flashinfer 0.6.13 and 0.6.14 (the
version currently pinned here). PR #3973 (merged into v0.6.15) is suggested
by a flashinfer maintainer as a possible fix, though not yet confirmed against
#3971 specifically. We have live py-spy/gdb captures from a production
dynamo-vllm Kimi-K2.6 hang (dflash + flashinfer_trtllm MoE, GB300, flashinfer
0.6.14) showing TP worker processes stuck mid-cuLaunchKernelEx inside
flashinfer::trtllm_fp4_block_scale_moe for the full ~9 minute hang window,
consistent with (but not proven identical to) #3971's failure mode.

Bumps flashinfer-commit v0.6.14 -> v0.6.15 in both CUDA variants of the
vllm-commit 752a3a5 (v0.25.1) matrix entry, matching the exact vLLM ref
currently deployed.

Also rebases the CoreWeave GLM-5.2 DeepSeekV3-no-group-routing patch
(originally against v0.6.14) onto v0.6.15:
- Its second hunk (csrc/trtllm_fused_moe_runner.cu, setting
  routingData.mSumEpsilon = 1e-20f as a 0/0-NaN guard) no longer applies:
  upstream's own v0.6.14->v0.6.15 diff independently adds the identical line
  at the same call site with an equivalent explanatory comment. Verified by
  diffing the real v0.6.15 source directly - the fix is now redundant
  upstream, not just textually conflicting. Dropped from the patch.
- Its first hunk (RoutingCustomPolicy.cuh, warp-per-token applyWithScores)
  still applies cleanly against v0.6.15 (verified via `git apply` against the
  real v0.6.15 source, 1-line offset). NOT verified: whether it's still
  necessary/correct given upstream also generalized the surrounding
  postprocess-dispatch architecture in this version (new kNeedsAux /
  applyWithAux / PolicyPairNeedsAux traits). Needs review from the patch's
  owner before this merges - flagged in the patch file itself.

Draft PR: this is a test build to gather evidence on the GB300 hang, not a
verified fix. Do not merge without confirming (a) the hang no longer
reproduces under load, over multiple runs, and (b) the GLM-5.2 patch rebase
above is correct.

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
@github-actions

Copy link
Copy Markdown

@nehashah-eng Build complete, success: https://github.com/coreweave/ml-containers/actions/runs/29860686764
Image: ghcr.io/coreweave/ml-containers/vllm-tensorizer:neha-flashinfer-0.6.15-gb300-moe-hang-test-f28b81b-752a3a504485790a2e8491cacbb35c137339ad34-cuda13.2.1-ubuntu24.04

@github-actions

Copy link
Copy Markdown

@nehashah-eng Build complete, success: https://github.com/coreweave/ml-containers/actions/runs/29860686764
Image: ghcr.io/coreweave/ml-containers/vllm-tensorizer:neha-flashinfer-0.6.15-gb300-moe-hang-test-f28b81b-752a3a504485790a2e8491cacbb35c137339ad34-cuda12.9.1-ubuntu24.04

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant