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(); }