我正在努力让Mark Harris减少CUDA的减少技巧＃5 。
减少＃5通过应用最后的扭曲展开来改善先前的减少＃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);
我不明白为什么没有__syncthreads()线之间一个和线乙（和下一线之间为好）。
我的问题：是否有可能在同一线程中，一个线程先执行B行，而另一个线程先执行A行？（似乎不可能，请任何人确认并详细说明）
在同一经纱中，一个线程是否可以在另一线程执行线A之前执行B行？
在撰写本文时（大约10年前），不可能发生这种情况，因为保证了翘曲可以在锁定步骤中执行。请注意，需要声明有问题的内存，volatile以防止编译器优化在Fermi和较新的GPU的缩减步骤之间缓存结果。在不需要的原始Tesla架构上。
但是，执行扭曲级操作的最新方式已经改变，并且这种类型的设计模式在某些最新的体系结构上可能是不安全的。取而代之的是，您应该使用扭曲级别原语进行缩减，而不是隐式扭曲同步。有关更多信息，请参见此博客文章。
本文收集自互联网，转载请注明来源。
如有侵权，请联系[email protected] 删除。
我来说两句