diff --git a/.github/configurations/vllm-tensorizer.yml b/.github/configurations/vllm-tensorizer.yml index 72268f5..2215ea2 100644 --- a/.github/configurations/vllm-tensorizer.yml +++ b/.github/configurations/vllm-tensorizer.yml @@ -1,13 +1,13 @@ include: # renovate-release: datasource=github-releases depName=vllm-project/vllm currentValue=v0.25.1 versioning=pep440 - vllm-commit: '752a3a504485790a2e8491cacbb35c137339ad34' - flashinfer-commit: 'v0.6.14' + flashinfer-commit: 'nightly-v0.6.15-20260724' 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' # renovate-release: datasource=github-releases depName=vllm-project/vllm currentValue=v0.25.1 versioning=pep440 - vllm-commit: '752a3a504485790a2e8491cacbb35c137339ad34' - flashinfer-commit: 'v0.6.14' + flashinfer-commit: 'nightly-v0.6.15-20260724' 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..257cd0f 100644 --- a/vllm-tensorizer/Dockerfile +++ b/vllm-tensorizer/Dockerfile @@ -97,9 +97,6 @@ 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 -RUN git -C flashinfer apply -v --stat --apply /git/patch && \ - rm /git/patch FROM alpine/git:2.36.3 AS lmcache-downloader diff --git a/vllm-tensorizer/flashinfer-v0.6.14-deepseekv3-no-group-routing.patch b/vllm-tensorizer/flashinfer-v0.6.14-deepseekv3-no-group-routing.patch deleted file mode 100644 index b169a1d..0000000 --- a/vllm-tensorizer/flashinfer-v0.6.14-deepseekv3-no-group-routing.patch +++ /dev/null @@ -1,98 +0,0 @@ -From 0000000000000000000000000000000000000000 Mon Sep 17 00:00:00 2001 -From: Ace Eldeib -Date: Thu, 2 Jul 2026 01:10:00 -0400 -Subject: [PATCH] Fix FlashInfer TRTLLM DeepSeek no-group routing weights - -GLM-5.2 uses FlashInfer TRTLLM NvFP4 MoE through the DeepSeekV3 no-group -routing path. That path selects experts with `sigmoid(router_logit) + -routing_bias`, then normalizes the selected experts with the unbiased sigmoid -scores. - -The block-per-token routing kernel already preserves that separation: it writes -the biased selection key to shared memory and writes the unbiased sigmoid into -an auxiliary shared-memory array consumed by `ScaledSumNormalizePostprocess`. -The warp-per-token routing path did not. It selected top-k using the biased -score, then tried to recover the unbiased sigmoid as `selected_score - bias`. - -That reconstruction is lossy for GLM-5.2. The model can produce tiny sigmoid -scores around 1e-8 to 1e-7 with routing correction biases around 12. In fp32, -adding the tiny sigmoid to the large bias rounds back to the bias, so -subtracting the bias later yields zero. When every selected expert in a row -collapses that way, `ScaledSumNormalizePostprocess` computes `0 / 0`, produces -NaN expert weights, and GLM decode can collapse into repeated punctuation. - -Keep expert selection unchanged, but recompute the selected experts' unbiased -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 ---- - .../trtllm/fused_moe/RoutingCustomPolicy.cuh | 30 +++++++++++++++++-- - csrc/trtllm_fused_moe_runner.cu | 1 + - 2 files changed, 29 insertions(+), 2 deletions(-) - -diff --git a/include/flashinfer/trtllm/fused_moe/RoutingCustomPolicy.cuh b/include/flashinfer/trtllm/fused_moe/RoutingCustomPolicy.cuh -index 934dcbf71..5f1f6a6ae 100644 ---- a/include/flashinfer/trtllm/fused_moe/RoutingCustomPolicy.cuh -+++ b/include/flashinfer/trtllm/fused_moe/RoutingCustomPolicy.cuh -@@ -434,6 +434,24 @@ struct ScaledSumNormalizePostprocess { - } - } - -+ template -+ __forceinline__ __device__ static void applyWithScores( -+ cg::thread_block_tile const& warp, DataType (&warpTopKScore)[K], -+ int32_t const (&warpTopKExpertIdx)[K], int32_t laneIdx, int32_t topK, -+ InputType const* ptrScores, ParamsT const& params) { -+ float sigmoidScore = 0.f; -+ if (laneIdx < topK) { -+ int32_t const expertIdx = warpTopKExpertIdx[laneIdx]; -+ sigmoidScore = sigmoid_accurate(static_cast(ptrScores[expertIdx])); -+ } -+ -+ float sum = cg::reduce(warp, sigmoidScore, cg::plus()); -+ if (laneIdx < topK) { -+ warpTopKScore[laneIdx] = -+ static_cast(sigmoidScore * params.routeScale / (sum + params.sumEpsilon)); -+ } -+ } -+ - /// Block-per-token variant: read un-biased sigmoid directly from smemAux - /// (written by SigmoidBiasPreprocess::applyToSmem) instead of reloading bias - /// 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); -+ if constexpr (std::is_same_v) { -+ PostprocessPolicy_::template applyWithScores( -+ warp, warpTopKScore, warpTopKExpertIdx, laneIdx, topK, ptrScores, -+ params.mExpertSelectParams.mPostprocessParams); -+ } else { -+ PostprocessPolicy_::apply(warp, warpTopKScore, warpTopKExpertIdx, laneIdx, topK, -+ params.mExpertSelectParams.mPostprocessParams); -+ } - } - }; - -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