From 28a2472f9588bb5ff429bba4f7a33062daffcaf5 Mon Sep 17 00:00:00 2001 From: Ishaan <101312037+guptaishaan@users.noreply.github.com> Date: Sun, 2 Aug 2026 19:59:42 -0700 Subject: [PATCH] [DeepSeek-V4] Fix nvcc 13 crash building the topk_v2 kernel (#32910) Co-authored-by: guptaishaan Co-authored-by: Xiaoyu Zhang <1182563586@qq.com> --- python/sglang/kernels/jit/csrc/deepseek_v4/topk_v2.cuh | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/python/sglang/kernels/jit/csrc/deepseek_v4/topk_v2.cuh b/python/sglang/kernels/jit/csrc/deepseek_v4/topk_v2.cuh index 2fb2b533a..88e6e9ae4 100644 --- a/python/sglang/kernels/jit/csrc/deepseek_v4/topk_v2.cuh +++ b/python/sglang/kernels/jit/csrc/deepseek_v4/topk_v2.cuh @@ -220,8 +220,13 @@ CLUSTER_TOPK_KERNEL void topk_small_batch_kernel(const __grid_constant__ TopKLau if (blockIdx.y == worker_rank) Streaming::forward(problem, &smem); } else { auto cluster = cooperative_groups::this_cluster(); - problem.out = cluster.map_shared_rank(topk_indices, worker_rank); - Cluster::forward(problem, &smem); // write to peer's output shared memory + // The mapped alias stays in a copy: the elected rank reads the very same + // bytes back through `topk_indices` below, and letting a shared::cluster + // address reach the `problem.out` that problem_transform loads makes cicc + // segfault on CUDA 13.x (issue #32830). + auto peer_problem = problem; + peer_problem.out = cluster.map_shared_rank(topk_indices, worker_rank); + Cluster::forward(peer_problem, &smem); // write to peer's output shared memory cluster.sync(); }