Warp内循环展开 & 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];
    }

上面的代码时在归约到剩下小于等于32个线程的情况会采用的展开方式，可以了解到使用的方法是Warp内部的循环展开。
之所以这样优化一是因为Warp内部具有隐性同步的特性（后面会详细叙述），二是线程小于32时会在一个Warp中做syncThreads()，显然具有不高效的问题
但是初看代码，好像存在Race Condition，并且还有重复计算的情况。举例vmem[1]会被多次累加，或者在使用vmem[1]时，vmem[1]的相关计算都完成了么。
在了解到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线程早已经把数据写完了；
3. 另外还需注意一点，当设置的BlockSize不是32的倍数时，发现最后的余数部分时最先执行的，也算是Warp调度的一个特点了

Warp调度顺序的实验代码：

__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(64, 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);
}