fix: 修复 online softmax warp_shuffle kernel 中 WarpReduceOnline 函数的计算边界问题 - #48
Open
kleinblue4 wants to merge 1 commit into
Open
kleinblue4 wants to merge 1 commit into
kleinblue4 wants to merge 1 commit into
Conversation
Author
|
有人来审一下 PR 么 |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
问题
原文“模块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,该误差符合正常的数据偏移