Bug起源
来自于Nvidia 的cub库的issue#179。当使用非2的次幂个逻辑warp大小时，归约的数据不对，具体可以参见上面的issue。我在使用时也遇见了类似问题。最后经过我的debug，找到了问题所在，并提交了pull request。
Bug 产生的原因
Bug产生的原因是由于lane_id的计算结果错误，当逻辑线程束的大小不是2的次幂时，将会使用shared memory进行归约，而WarpReduceSmem的类初始化计算lane_id时，计算错误。这个错误只会出现在当逻辑线程束跨越两个物理线程束时。
/// Constructor
__device__ __forceinline__ WarpReduceSmem(
TempStorage &temp_storage)
:
temp_storage(temp_storage.Alias()),
lane_id(IS_ARCH_WARP ?
LaneId() :
LaneId() % LOGICAL_WARP_THREADS),// 这里计算错误
member_mask((0xffffffff >> (32 - LOGICAL_WARP_THREADS)) <<
((IS_ARCH_WARP || !IS_POW_OF_TWO ) ?
0 : // arch-width and non-power-of-two subwarps cannot be tiled with the arch-warp
((LaneId() / LOGICAL_WARP_THREADS) * LOGICAL_WARP_THREADS)))
{}
为了明白说明，举个例子。如LOGICAL_WARP_THREADS=7，每个block的线程数为5。那么对于thereadIdx.x=28时为第五个逻辑线程束的开始，那么此逻辑线程束会跨越第一个物理线程束和第二个物理线程束。结果如下：
tid=28-31正确的lane_id
tid=32-34正确的lane_id
从上面可以看出，当到了第二个物理线程束时，它的lane_id又从0开始计数了，这是不对的。所以会出现错误的结果。它和正确的结果对比如下：
| threadId | 28 | 29 | 30 | 31 | 32 | 33 | 34 |
|---|---|---|---|---|---|---|---|
| lane ID(incorrect) | 0 | 1 | 2 | 3 | 0 | 1 | 2 |
| lane ID(correct) | 0 | 1 | 2 | 3 | 4 | 5 | 6 |
当lane_id在后面的进行的shred memory方式的归约时，使用它取用的共享内存的数组将会是错误的：
// Update input if peer_addend is in range
if ((ALL_LANES_VALID && IS_POW_OF_TWO) || ((lane_id + OFFSET) < valid_items))
{
// 下面的temp_storage.reduce[lane_id + OFFSET]由于lane_id错误，将导致数据错误
T peer_addend = ThreadLoad<LOAD_VOLATILE>(&temp_storage.reduce[lane_id + OFFSET]);
input = reduction_op(input, peer_addend);
}
Bug的解决
在知道了bug的出处后，解决就变得容易了，我们应当根据线程号而不是LaneId()号来计算逻辑线程束的lane_id，因此，最终的修改是：
- 修改1
修改1
这样修改虽然看似可以，但是这也只是在线程块为1维或者在线程块y和z维度的大小是物理线程束大小的整数倍时有效。
- 修改2
修改2
因此，应当某个线程总的线程号，以此来计算逻辑线程束下的lane_id
经过这样修改后，计算结果正确。
结论
- bug产生原因是由于
lane_id计算错误 - 一步一步仔细debug是找到bug的方法
