Skip to content

fix: 修复 online softmax warp_shuffle kernel 中 WarpReduceOnline 函数的计算边界问题 - #48

Open
kleinblue4 wants to merge 1 commit into
caomaolufei:mainfrom
kleinblue4:fix/Online_Softmax_Warp_Shuffle_Kernel
Open

kleinblue4 wants to merge 1 commit into
caomaolufei:mainfrom
kleinblue4:fix/Online_Softmax_Warp_Shuffle_Kernel

Conversation

@kleinblue4

Copy link
Copy Markdown

问题

原文“模块2-CUDA编程与算子优化/5.2-CUDA Online Softmax实现/4.3 Kernel实现”中,warpReduceOnline 函数默认设置 offset=16,但在实际执行过程中,softmax 的计算结果为 NaN

原因

在4.3 Kernel实现代码中的第二级 warp 规约,其执行规约的线程数可能不为32,即num_warps 不为32,假设 blockDim.x = 256,num_warps = 8,对于 thread 0-7,local_m = warp_m[lane],但对于 thread 8-31,local_m = -INFINITY,那么考虑在 warpReduceOnline 函数中,当 offset = 16时,thread 9 获取到 thread 25 的 m2 = -INFINITY,此时 m = m2 = -INFINITY,这会导致后续计算 m - m_new = NaN,从而 expf(Nan) = NaN,使得最终计算的 sum = NaN

修复

在 warpReduceOnline 函数中添加了 Warp_Threads 变量,用于指示当前 warp 中实际参与计算的线程数量,从而确保计算过程中不会涉及到无关的线程

测试

通过将修改后代码的计算结果与 libtorch 中的 torch::softmax 计算结果进行对比,两者的最大绝对误差为 3.027e-09,该误差符合正常的数据偏移

@kleinblue4 kleinblue4 changed the title 修复 online softmax warp_shuffle kernel 中 WarpReduceOnline 的bug fix: 修复 online softmax warp_shuffle kernel 中 WarpReduceOnline 的bug Sep 21, 2026
@kleinblue4 kleinblue4 changed the title fix: 修复 online softmax warp_shuffle kernel 中 WarpReduceOnline 的bug fix: 修复 online softmax warp_shuffle kernel 中 WarpReduceOnline 的 Sep 21, 2026
@kleinblue4 kleinblue4 changed the title fix: 修复 online softmax warp_shuffle kernel 中 WarpReduceOnline 的 fix: 修复 online softmax warp_shuffle kernel 中 WarpReduceOnline 函数的计算边界问题 Sep 21, 2026
@kleinblue4

Copy link
Copy Markdown
Author

有人来审一下 PR 么

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant