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

参考documentation、atomicInc 这样做：

以下：

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

计算：

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

并将结果存储回count[0]

（如果您不确定? 运算符的作用，请查看here）

由于您已经传递了n = 1，并且count[0] 从 1 开始，atomicInc 实际上从未将变量 count[0] 增加超过 1。

如果您希望看到它的增量超过 1，请为 n 传递一个更大的值。

变量n 实际上充当递增过程的“翻转值”。当要递增的变量实际达到n的值时，下一个atomicInc会将其重置为零。

虽然您没有问过这个问题，但您可能会问：“如果我达到了翻转值，为什么我永远看不到零值？”

要回答这个问题，您必须记住所有 4 个线程都在同步执行。它们都执行atomicInc指令，然后再执行后续的打印语句。

因此，我们有一个变量 count[0]，它从 1 开始。

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

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

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

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

如 cmets 中的 @talonmies 所示，如果您将 atomicInc 换成 atomicAdd：

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

你会得到你可能期望的结果。