From c9c96b63f765db4ac7c88e824a755851f1049e30 Mon Sep 17 00:00:00 2001 From: EanWang211123 Date: Mon, 27 Jul 2026 11:07:16 +0800 Subject: [PATCH 1/2] fix compact graph prefill plan race Signed-off-by: EanWang211123 --- python/sglang/kernels/jit/csrc/deepseek_v4/c_plan.cuh | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/python/sglang/kernels/jit/csrc/deepseek_v4/c_plan.cuh b/python/sglang/kernels/jit/csrc/deepseek_v4/c_plan.cuh index c86b92135437..eb3cf86c1abe 100644 --- a/python/sglang/kernels/jit/csrc/deepseek_v4/c_plan.cuh +++ b/python/sglang/kernels/jit/csrc/deepseek_v4/c_plan.cuh @@ -154,6 +154,10 @@ __global__ __launch_bounds__(1024, 1) // warp_max[tx] = 0; warp_min[tx] = 0xFFFFFFFFu; } + // Warp 0 initializes all warp_min slots; each warp then writes its own. Without + // this barrier, init can clobber a reduction result and make ragged extend look + // uniform, wrongly selecting the MTP fast path and overrunning ragged_id. + __syncthreads(); // === Stage B: min/max(extend_len) for MTP-uniform detection === // For min, treat threads outside `batch_size` as +inf so they don't pull the min down. From 58124fb983136b903108c8ef1305c135eb7d376f Mon Sep 17 00:00:00 2001 From: EanWang211123 Date: Mon, 10 Aug 2026 20:14:31 +0800 Subject: [PATCH 2/2] [update] delete the useless initation Signed-off-by: EanWang211123 --- python/sglang/kernels/jit/csrc/deepseek_v4/c_plan.cuh | 9 --------- 1 file changed, 9 deletions(-) diff --git a/python/sglang/kernels/jit/csrc/deepseek_v4/c_plan.cuh b/python/sglang/kernels/jit/csrc/deepseek_v4/c_plan.cuh index eb3cf86c1abe..5e9164717bb7 100644 --- a/python/sglang/kernels/jit/csrc/deepseek_v4/c_plan.cuh +++ b/python/sglang/kernels/jit/csrc/deepseek_v4/c_plan.cuh @@ -150,15 +150,6 @@ __global__ __launch_bounds__(1024, 1) // counter_c = 0; counter_w = 0; } - if (tx < kNumWarps) { - warp_max[tx] = 0; - warp_min[tx] = 0xFFFFFFFFu; - } - // Warp 0 initializes all warp_min slots; each warp then writes its own. Without - // this barrier, init can clobber a reduction result and make ragged extend look - // uniform, wrongly selecting the MTP fast path and overrunning ragged_id. - __syncthreads(); - // === Stage B: min/max(extend_len) for MTP-uniform detection === // For min, treat threads outside `batch_size` as +inf so they don't pull the min down. const uint32_t e_for_max = static_cast(extend_len);