CUDA的atomicAdd示例无法产生正确的输出。

Question

以下代码旨在将100个浮点数元素的数组递增1十次。在输出中，我希望每个元素都是10.0f值的100元素数组。但是实际上，我得到了随机值。您能指出我的错误吗？

__global__  void testAdd(float *a)
{
    float temp;
    for (int i = 0; i < 100 ; i++)
    {
        a[i] = atomicAdd(&a[i], 1.0f);
    }
}
void cuTestAtomicAdd(float *a)
{
    testAdd<<<1, 10>>>(a);
}

我的目标是理解原子操作的工作方式，以便将其应用于其他地方。

- Gumbly jr.

---

我向您保证，变量'a'已经正确分配了100个元素的空间，并且所有元素都已经被初始化为0。 - Gumbly jr.

为了避免竞态条件，在C++中，每次读写内存时我会使用互斥锁。在CUDA中有相应的等价物吗？ - Gumbly jr.

此外，我正在GTX 780 Ti上运行此程序，我认为它应该具有足够高的计算能力来支持原子操作。我正在使用Visual Studio 2010。我已经更改了所有.cpp和.cu文件（Device\Code Generation）的属性，从之前的compute_10、sm_10更改为compute_35、sm_35。 - Gumbly jr.

Answer 1

这不是我们执行atomicAdd操作的方法。

应该像这样执行：

atomicAdd(&a[i], 1.0f);

并且被修改的变量 (a[i]) 将会被更新。

原子函数的返回值通常是在原子更新之前，该变量中的旧值。

因此执行以下操作：

a[i] = atomicAdd(&a[i], 1.0f);

将更新变量a[i]，然后（非原子地）将旧值分配给变量a[i]。这几乎肯定不是您想要的。

阅读文档：

该函数返回旧值。

下面的完整代码演示了正确的用法：

#include <iostream>

__global__  void testAdd(float *a)
{
    for (int i = 0; i < 100 ; i++)
    {
        atomicAdd(&a[i], 1.0f);
    }
}
void cuTestAtomicAdd(float *a)
{
    testAdd<<<1, 10>>>(a);
}

int main(){

  float *d_data, *h_data;
  h_data=(float *) malloc(100*sizeof(float));
  cudaMalloc((void **)&d_data, 100*sizeof(float));
  cudaMemset(d_data, 0, 100*sizeof(float));
  cuTestAtomicAdd(d_data);
  cudaMemcpy(h_data, d_data, 100*sizeof(float), cudaMemcpyDeviceToHost);
  for (int i = 0; i < 100; i++)
    if (h_data[i] != 10.0f) {printf("mismatch at %d, was %f, should be %f\n", i, h_data[i], 10.0f); return 1;}
  printf("Success\n");
  return 0;
}

- Robert Crovella