atomicInc()不起作用

Question

我使用atomicInc()尝试过以下程序.

__global__ void ker(int *count)
{
    int n=1;
    int x = atomicInc ((unsigned int *)&count[0],n);
    CUPRINTF("In kernel count is %d\n",count[0]);
}

int main()
{
    int hitCount[1];
    int *hitCount_d;

    hitCount[0]=1;
    cudaMalloc((void **)&hitCount_d,1*sizeof(int));

    cudaMemcpy(&hitCount_d[0],&hitCount[0],1*sizeof(int),cudaMemcpyHostToDevice);

    ker<<<1,4>>>(hitCount_d);

    cudaMemcpy(&hitCount[0],&hitCount_d[0],1*sizeof(int),cudaMemcpyDeviceToHost);

    printf("count is %d\n",hitCount[0]);
  return 0;
}

输出是:

In kernel count is 1
In kernel count is 1
In kernel count is 1
In kernel count is 1

count is 1

我不明白为什么它没有递增.谁能帮忙

Answer 1

参考文档,atomicInc这样做:

对于以下内容:

atomicInc ((unsigned int *)&count[0],n);

计算:

((count[0] >= n) ? 0 : (count[0]+1))

并将结果存回 count[0]

(如果您不确定?操作员的操作,请查看此处)

由于您已经传递n= 1,并且count[0]从1 开始,因此 atomicInc实际上不会将变量count[0]增加到1以上.

如果你想看到它增加超过1,则传递一个更大的值n.

该变量n实际上充当递增过程的"翻转值".当要递增的变量实际达到值时n,下一个atomicInc将其重置为零.

虽然你没有问过这个问题,但你可能会问,"如果我达到翻转价值,为什么我从未看到零值？"

要回答这个问题,您必须记住所有4个线程都是以锁步方式执行的.atomicInc在执行后续print语句之前,所有4个都执行该指令.

因此,我们的变量count[0]从1开始.

1. 执行原子的第一个线程将其重置为零.
2. 下一个线程将其递增为1.
3. 第三个线程将其重置为零.
4. 第四个也是最后一个线程将其递增为1.

然后所有4个线程打印出该值.

作为另一个实验,尝试启动5个线程而不是4个线程,看看您是否可以预测打印出来的值.

ker<<<1,5>>>(hitCount_d);

正如@talonmies在评论中指出的那样,如果你换atomicInc了一个atomicAdd:

int x = atomicAdd ((unsigned int *)&count[0],n);

你会得到你可能期待的结果.