我正在努力做到在Nvidia Reduction上看到的所有优化.我已经实现了前四个部分,但我在第22个幻灯片中遇到了第5部分.
我无法理解为什么提供的代码可以在没有任何syncthreads()的情况下工作的原因.线程可以访问输出中的相同内存位置.
此外,幻灯片表明,如果变量未设置为volatile,则代码将不起作用.在这方面如何变得不稳定？如果我不想调用内核,那么编程它的最佳方法是什么？
我也把这些代码放在这里供参考.
__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];
}
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);
在此先感谢您的帮助.如果需要更多信息,请评论.