树约简中共享内存的几个问题

Question

在NVIDIA 优化CUDA的并行约简第22页，

__device__ void warpReduce(volatile int* sdata, int tid) {
sdata[tid] += sdata[tid + 32];
sdata[tid] += sdata[tid + 16];
sdata[tid] += sdata[tid + 8];
sdata[tid] += sdata[tid + 4];
sdata[tid] += sdata[tid + 2];
sdata[tid] += sdata[tid + 1];
}

// later…
for (unsigned int s=blockDim.x/2; s>32; s>>=1) {
if (tid < s)
    sdata[tid] += sdata[tid + s];
__syncthreads();
}
if (tid < 32) warpReduce(sdata, tid);

我不确定在函数warpReduce(...)：sdata[tid] += sdata[tid + 1];的最后一行中CUDA是如何处理操作的。如果翘曲从0到31有线程ID tid，那么tid 0和1在此时冲突吗？我的意思是，数据自动化系统将如何选择sdata[0] += sdata[0 + 1];和sdata[1] += sdata[1 + 1];的顺序？如果它首先运行sdata[0] += sdata[0 + 1];，则结果是正确的。或者，数据自动化系统做些什么来防止这种混乱/冲突？

Answer 1

那些幻灯片已经过时了。使用共享内存在一次翘曲范围内减少还不是最先进的，因为翘曲洗牌指令是在CUDA-3.0 (IIRC)中发明的。就连这些都已经过时了。

现在，您将使用经纱缩减功能 (如__reduce_add_sync )。

至于上面的模式为什么会起作用:在Volta体系结构出现之前，翘曲中的线程是以锁定的方式工作的。他们都会首先加载值，然后存储它。因此，没有必要进行同步。这种情况已经改变了。请注意，这来自Nvidia的伏尔塔调谐指南

假定读写的应用程序对处于同一翘曲中的其他线程隐式可见，需要在通过全局内存或共享内存在线程之间交换数据的步骤之间插入新的__syncwarp() warp宽障碍同步指令。假定代码是在锁步中执行的，或者是从单独的线程中读取/写的，则在没有同步的情况下，在翘曲中可以看到代码是无效的。

您可以在Github存储库中找到更新的示例代码：https://github.com/NVIDIA/cuda-samples。

也许还有一个新版本的幻灯片，但我不知道在哪里。