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开始计数了，这是不对的。所以会出现错误的结果。它和正确的结果对比如下：

当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

这样修改虽然看似可以，但是这也只是在线程块为1维或者在线程块y和z维度的大小是物理线程束大小的整数倍时有效。

1. 修改2

修改2

因此，应当某个线程总的线程号，以此来计算逻辑线程束下的lane_id

经过这样修改后，计算结果正确。

结论

1. bug产生原因是由于lane_id计算错误
2. 一步一步仔细debug是找到bug的方法