14.5 基于线程束洗牌实现

流处理器簇（SM）3.x新增加了线程束洗牌指令（参见8.6.1节的描述），让线程能够在寄存器之间交换数据而不用通过共享内存中转。_shf1()内置函数能够将某个线程的寄存器数据直接向同线程束中其他所有线程进行广播。就如代码清单14-4所展示的，代替线程块大小的分块和共享内存，我们可以使用尺寸为32的数据分块（对应线程束大小），然后在线程束内部通过广播将个体的数据在线程之间进行读取。

有趣的是，这种策略相对于共享内存实现方式降低了 $25\%$ 的性能（每秒340亿个相对于452亿个个体相互作用力计算）。线程洗牌指令的开销几乎等同共享内存的访问，而且计算的执行是按线程束（32线程）尺寸分块的而不是线程块的尺寸。看起来使用线程洗牌指令代替读写共享内存获得的效果更好，而不只是代替单纯的读共享内存。线程洗牌应该只被用于当内核需要共享内存做其他事的时候。

代码清单14-7 ComputeNBodyGravitation_Shuffle内核

global void
ComputeNBodyGravitation_Shuffle(
    float *force,
    float *posMass,
    float softeningSquared,
    size_t N)
{
    const int laneid = threadIdx.x & 31;
    for (int i = blockIdx.x*blockDim.x + threadIdx.x;
        i < N;
            i += blockDim.x*gridDim.x)
        {
            float acc[3] = {0};
            float4 myPosMass = ((float4 *) posMass)[i];
        }
        for (int j = 0; j < N; j++) {
            float4 shufSrcPosMass = ((float4 *) posMass)[j+laneid];
        }
    for (int k = 0; k < 32; k++) {
        float fx, fy, fz;
        float4 shufDstPosMass;
        shufDstPosMass.x = _shfl(shufSrcPosMass.x, k);
        shufDstPosMass.y = _shfl(shufSrcPosMass.y, k);
        shufDstPosMass.z = _shfl(shufSrcPosMass.z, k);
        shufDstPosMass.w = _shfl(shufSrcPosMass.w, k);
        bodyBodyInteraction(
            &fx, &fy, &fz,
            myPosMass.x, myPosMass.y, myPosMass.z,
            shufDstPosMass.x,
            shufDstPosMass.y,
            shufDstPosMass.z,
            shufDstPosMass.w,
            softeningSquared);
            acc[0] += fx;
            acc[1] += fy;
            acc[2] += fz;
        }
    }
    force[3*i+0] = acc[0];
    force[3*i+1] = acc[1];
    force[3*i+2] = acc[2];
}