问题描述
我正在努力让Mark Harris着手reduction in CUDA中的还原技术＃5。
Reduction＃5通过应用最后的翘曲展开来改善先前的Reduction＃4。
幻灯片21提到：“我们不需要__syncthreads()”，这是我不理解的部分。
以下是具有主要逻辑的代码：
__device__ void warpReduce(volatile int* sdata,int tid) {
sdata[tid] += sdata[tid + 32]; // line A
sdata[tid] += sdata[tid + 16]; // line B
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);
我不明白为什么在 A 行和 B 行之间以及接下来的行之间没有__syncthreads()。
我的问题：是否可能在同一个扭曲内，一个线程执行 B 行，而另一个线程执行 A 行？ （似乎不可能，请任何人确认并详细说明）
解决方法
是否有可能在同一线程中，一个线程先执行B行，而另一个线程先执行A行？
在撰写本文时（大约10年前），不可能发生这种情况，因为保证了翘曲可以在锁定步骤中执行。请注意，需要声明有问题的内存volatile，以防止编译器优化在Fermi和较新的GPU的简化步骤之间缓存结果。不需要原来的Tesla架构。
但是，执行扭曲级操作的最新方式已经改变，并且这种类型的设计模式在某些最新的体系结构上可能是不安全的。取而代之的是，您应该更喜欢使用扭曲级别的原语进行缩减，而不是隐式扭曲同步。有关更多信息，请参见this blog post。