目的:想求一个数组 A[100] 中的所有大于0的数,并且将他们的值分别保存到数组B中。
假设一个计数器k=0;启动100个线程,然后在kernel函数中
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if(A[tid]>0)
{
B[k]=A[tid];
k++;
}
可是这样计算k的值无法同步,导致计算出的B数组不正确。
能不能将if 函数里面的两行代码设置成最小操作,类似于atomic function?
这样保证k在加1的时候,B的索引能够正确。
楼主您好,您的设计思路错误:
对于一个维护了尾部位置的写入缓冲区来说,此2步没必要被原子性的完成。
而只需要将尾部索引的+1和返回旧值原子完成即可:
int tail = atomicAdd(&k, 1);
B[tail] = …;
这样是非常安全的。
以及,如果您需要的是循环缓冲区的话,请改用atomicInc:
int tail = atomicInc(&k, BUFFER_LENGTH_IN_ELEMENTS); //如果+1后超出最大尾部索引,将改成0位置。
B[tail] = …;
这也是非常安全的。
感谢您莅临CUDAZone China。
感谢你的指点,果然得到了正确的结果。
我用的是第一种方法,貌似第二个方法的标志BUFFER_LENGTH_IN_ELEMENTS在程序中无法找到定义。