minke.yu/sglang
1
0
Fork 0
Code Issues Pull Requests Actions 12 Packages Projects Releases Wiki Activity

[DeepSeek-V4] Fix nvcc 13 crash building the topk_v2 kernel (#32910)

Browse Source 
Co-authored-by: guptaishaan <guptaishaan@users.noreply.github.com>
Co-authored-by: Xiaoyu Zhang <1182563586@qq.com>
This commit is contained in:
Ishaan
2026-08-03 10:59:42 +08:00
committed by GitHub
co-authored by guptaishaan Xiaoyu Zhang
parent 4bc593fdc8
commit 28a2472f95
1 changed files with 7 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files
python/sglang/kernels/jit/csrc/deepseek_v4/topk_v2.cuh
+7 -2
View File
@@ -220,8 +220,13 @@ CLUSTER_TOPK_KERNEL void topk_small_batch_kernel(const __grid_constant__ TopKLau
if (blockIdx.y == worker_rank) Streaming::forward<kPDL>(problem, &smem);
} else {
auto cluster = cooperative_groups::this_cluster();
problem.out = cluster.map_shared_rank(topk_indices, worker_rank);
Cluster::forward<kPDL>(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<kPDL>(peer_problem, &smem); // write to peer's output shared memory
cluster.sync();
}