diff --git a/.github/configurations/vllm-tensorizer.yml b/.github/configurations/vllm-tensorizer.yml index d1a0eaa..a94f6cf 100644 --- a/.github/configurations/vllm-tensorizer.yml +++ b/.github/configurations/vllm-tensorizer.yml @@ -1,11 +1,11 @@ include: - vllm-commit: '752a3a504485790a2e8491cacbb35c137339ad34' - flashinfer-commit: 'v0.6.14' + flashinfer-commit: 'v0.6.15' lmcache-commit: 'v0.4.7' builder-base-image: 'ghcr.io/coreweave/ml-containers/torch:bc8c66e-nccl-cuda13.2.1-ubuntu24.04-nccl2.30.4-1-torch2.11.0-vision0.26.0-audio2.11.0-abi1' final-base-image: 'ghcr.io/coreweave/ml-containers/torch:bc8c66e-nccl-cuda13.2.1-ubuntu24.04-nccl2.30.4-1-torch2.11.0-vision0.26.0-audio2.11.0-abi1' - vllm-commit: '752a3a504485790a2e8491cacbb35c137339ad34' - flashinfer-commit: 'v0.6.14' + flashinfer-commit: 'v0.6.15' lmcache-commit: 'v0.4.7' builder-base-image: 'ghcr.io/coreweave/ml-containers/torch:bc8c66e-nccl-cuda12.9.1-ubuntu24.04-nccl2.30.4-1-torch2.11.0-vision0.26.0-audio2.11.0-abi1' final-base-image: 'ghcr.io/coreweave/ml-containers/torch:bc8c66e-nccl-cuda12.9.1-ubuntu24.04-nccl2.30.4-1-torch2.11.0-vision0.26.0-audio2.11.0-abi1' diff --git a/vllm-tensorizer/Dockerfile b/vllm-tensorizer/Dockerfile index c0154e0..669052f 100644 --- a/vllm-tensorizer/Dockerfile +++ b/vllm-tensorizer/Dockerfile @@ -97,7 +97,7 @@ RUN git clone --filter=tree:0 --no-single-branch --no-checkout \ git checkout "${FLASHINFER_COMMIT}" && \ git submodule update --init --recursive --jobs 8 \ --depth 1 --filter=tree:0 -COPY flashinfer-v0.6.14-deepseekv3-no-group-routing.patch /git/patch +COPY flashinfer-v0.6.15-deepseekv3-no-group-routing.patch /git/patch RUN git -C flashinfer apply -v --stat --apply /git/patch && \ rm /git/patch diff --git a/vllm-tensorizer/flashinfer-v0.6.14-deepseekv3-no-group-routing.patch b/vllm-tensorizer/flashinfer-v0.6.15-deepseekv3-no-group-routing.patch similarity index 70% rename from vllm-tensorizer/flashinfer-v0.6.14-deepseekv3-no-group-routing.patch rename to vllm-tensorizer/flashinfer-v0.6.15-deepseekv3-no-group-routing.patch index b169a1d..c54f227 100644 --- a/vllm-tensorizer/flashinfer-v0.6.14-deepseekv3-no-group-routing.patch +++ b/vllm-tensorizer/flashinfer-v0.6.15-deepseekv3-no-group-routing.patch @@ -26,15 +26,36 @@ sigmoid scores from the original logits before the warp-path normalization. This matches the existing block-per-token auxiliary path and the Python routing reference: selection uses `sigmoid + bias`; final weights use `sigmoid`. -Also set the same denominator epsilon already used by the adjacent MiniMax2 -path. The epsilon is only a final guard; the main correctness fix is preserving -the unbiased sigmoid for normalization. - Signed-off-by: Ace Eldeib + +--- REBASE NOTE (bump to flashinfer v0.6.15, this session) --- +Original patch (against v0.6.14) had two hunks: this one +(RoutingCustomPolicy.cuh, warp-per-token `applyWithScores`) and a second hunk +in csrc/trtllm_fused_moe_runner.cu that set `routingData.mSumEpsilon = 1e-20f;` +as a denominator-floor guard. + +The second hunk no longer applies against v0.6.15 and has been dropped from +this patch: upstream flashinfer's v0.6.14->v0.6.15 diff independently adds the +identical line (`routingData.mSumEpsilon = 1e-20f;`) at the same call site, +with an equivalent explanatory comment ("Floor the renorm denominator ... +1e-20f matches DeepSeek-V3's reference gate ... and the MiniMax2 branch") — +verified by fetching the real v0.6.15 source and diffing directly. That fix is +now redundant, not just textually conflicting. + +This first hunk (below) still applies cleanly against v0.6.15 (1-line offset, +verified via `git apply` against the real v0.6.15 source tree). NOT verified: +whether it is still *necessary* or *correct* given upstream also generalized +the surrounding postprocess-dispatch architecture in this version (new +`kNeedsAux` / `applyWithAux` / `PolicyPairNeedsAux` trait system replacing the +narrower block-per-token-only special-casing that existed in v0.6.14). This +needs a look from whoever owns this patch (Ace Eldeib) before merging — +flagging explicitly rather than guessing, since this is GLM-5.2 numerical +correctness territory, not the GB300/dflash hang this branch's version bump is +otherwise testing. --- - .../trtllm/fused_moe/RoutingCustomPolicy.cuh | 30 +++++++++++++++++-- - csrc/trtllm_fused_moe_runner.cu | 1 + - 2 files changed, 29 insertions(+), 2 deletions(-) + + .../trtllm/fused_moe/RoutingCustomPolicy.cuh | 18 ++++++++++++++++++ + 1 file changed, 18 insertions(+) diff --git a/include/flashinfer/trtllm/fused_moe/RoutingCustomPolicy.cuh b/include/flashinfer/trtllm/fused_moe/RoutingCustomPolicy.cuh index 934dcbf71..5f1f6a6ae 100644 @@ -43,7 +64,7 @@ index 934dcbf71..5f1f6a6ae 100644 @@ -434,6 +434,24 @@ struct ScaledSumNormalizePostprocess { } } - + + template + __forceinline__ __device__ static void applyWithScores( + cg::thread_block_tile const& warp, DataType (&warpTopKScore)[K], @@ -67,7 +88,7 @@ index 934dcbf71..5f1f6a6ae 100644 /// and subtracting. Saves one global memory read + one FMA per top-K lane. @@ -566,8 +584,14 @@ struct TopKExpertSelect { topk::reduceTopK(warp, warpTopKScore, warpTopKExpertIdx, score, idx, minScore, topK); - + // Apply postprocess (e.g. renormalize, softmax over top-K, scaled renormalize, ...) - PostprocessPolicy_::apply(warp, warpTopKScore, warpTopKExpertIdx, laneIdx, topK, - params.mExpertSelectParams.mPostprocessParams); @@ -81,18 +102,6 @@ index 934dcbf71..5f1f6a6ae 100644 + } } }; - -diff --git a/csrc/trtllm_fused_moe_runner.cu b/csrc/trtllm_fused_moe_runner.cu -index 6e7c97395..f2d949d51 100644 ---- a/csrc/trtllm_fused_moe_runner.cu -+++ b/csrc/trtllm_fused_moe_runner.cu -@@ -76,6 +76,7 @@ void Runner::run(void* routingLogits, void* routingBias, int32_t numTokens, int3 - routingData.mPtrRoutingBias = routingBias; - routingData.mDtypeBias = dtypeBias; - routingData.mRouteScale = routedScalingFactor; -+ routingData.mSumEpsilon = 1e-20f; - - routingData.mPtrScores = expertIds == nullptr ? routingLogits : nullptr; - routingData.mPtrTopKIds = expertIds; --- + +-- 2.50.0