Warp内部的硬件调度特性

// unrolling warp
if (tid < 32)
{
    volatile int *vmem = idata;
    vmem[tid] += vmem[tid + 32];
    vmem[tid] += vmem[tid + 16];
    vmem[tid] += vmem[tid +  8];
    vmem[tid] += vmem[tid +  4];
    vmem[tid] += vmem[tid +  2];
    vmem[tid] += vmem[tid +  1];
}

• 上面这段代码常用在Cuda累加运算的最后32个线程内的计算，利用循环展开来提升计算的并行度
• 初看这段代码应该会有疑惑：vmem的值没有多加么？最后都加到vmem[0]的值时最终的结果么？
• 比如：0号线程完成了tid+16了，此时16号线程实际上还没有运算玩，还有别的线程要往16号内存写数据
• 比如：0号线程0+32和1号线程1+31，这不是重复计算了么？
• 其实代码没有问题，而是需要对GPU的Warp内的硬件调度有深入理解才可以
• 原因1：Warp以32个线程为一组统一调度，这32个线程的每条指令都是同时执行的，比如32个线程同时执行“vmem[tid] += vmem[tid + 32];”，然后在同时执行“vmem[tid] += vmem[tid + 16];”。当知道此点时，我们即可以了解到第一步相当于32个线程同时将32到64的数据累加到了0到31的线程里；
• 原因2：Warp内部的内存事务是以16个线程为单位执行的，及halfWarp。而且是前16个线程先执行，再后16个线程执行。因此在第二部“vmem[tid] += vmem[tid + 16];”，前16个线程把16到31的数据写入0到15个线程内，之后才是16到31线程写后面的32到48的数据，但是显然我们已经不关系后面的16到31线程内部数据的结果了，他们只是作为Warp统一调度的最小执行单元必须跟着执行罢了。我们关注的0到15线程早已经把数据写完了；
• 另外还需注意一点，当设置的BlockSize不是32的倍数时，发现最后的余数部分时最先执行的，也算是Warp调度的一个特点了；看下面的测试代码，blockSize=48时，32至47号线程的结果时最先计算的

__global__ void checkIndex(void)
{
    printf("threadIdx:%d,blockIdx:%d\n", threadIdx.x, blockIdx.x);
}

int main(int argc, char** argv)
{
    // define grid and block structure
    dim3 block(48, 1);
    dim3 grid(3, 1);
    // check grid and block dimension from host side
    printf("grid.x %d grid.y %d grid.z %d\n", grid.x, grid.y, grid.z);
    printf("block.x %d block.y %d block.z %d\n", block.x, block.y, block.z);
    // check grid and block dimension from device side
    checkIndex << <grid, block >> > ();
    // reset device before you leave
    cudaDeviceReset();

    return(0);
}