CUDA atomicAdd 失败

Question

下面的 CUDA 内核应该对 3D 图像进行图像切片添加，即，您沿一维折叠 3D 体积并通过逐像素添加生成一张 2D 图像。 image_in 数据指针的大小为 128 * 128 * 128，它是使用函数 GetOutputBuffer() 从 ITK::Image 获得的。在阅读了 ITK 文档之后，我认为我们可以安全地假设数据指针指向图像数据的一段连续内存，没有填充。 image_out 只是大小为 128 * 128 的 2D 图像，也是由 ITK::Image 生成的。我包含有关图像的信息只是为了完整性，但问题更多的是关于 CUDA atomic 并且可能非常基本。该代码首先计算线程id，并将id投影到128 * 128的范围内，这意味着沿着我们执行加法的维度的同一行中的所有像素将具有相同的idx。然后使用这个 idx，atomicAdd 来更新 image_out。

__global__ void add_slices(int* image_in, int* image_out) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    int idx = tid % (128 * 128);
    int temp = image_in[tid];

    atomicAdd( &image_out[idx], temp );

}

我初始化image_out的方式是通过以下方式，有两种方式我尝试了类似的结果：

int* image_out = new int[128 * 128];
for (...) {
    /* assign image_out to zeros */
}

还有一个使用ITK接口的：

out_image->SetRegions(region2d);
out_image->Allocate();
out_image->FillBuffer(0);
// Obtain the data buffer
int* image_out = out_image->GetOutputBuffer();

然后我将 CUDA 设置如下：

unsigned int size_in = 128 * 128 * 128;
unsigned int size_out = 128 * 128;
int *dev_in;
int *dev_out;
cudaMalloc( (void**)&dev_in, size_in * sizeof(int) );
cudaMalloc( (void**)&dev_out, size_out * sizeof(int));
cudaMemcpy( dev_in, image_in, size_in * sizeof(int), cudaMemcpyHostToDevice );
add_slices<<<size_in/64, 64 >>>(dev_in, dev_out);
cudaMemcpy( image_out, dev_out, size_out * sizeof(int), cudaMemcpyDeviceToHost);

上面的代码有问题吗？我在这里寻求帮助的原因是因为上面的代码有时可能会产生正确的结果（我每运行 50 次代码，也许，我发誓我至少看到了两次正确的结果），而其余的当时只是产生了一些垃圾。问题是否来自 atomicAdd() 函数？一开始我的图像类型是double，CUDA不支持atomicAdd(double*, double)所以我使用Nvidia提供的代码如下

__device__ double atomicAdd(double* address, double val)
{
    unsigned long long int* address_as_ull =
                                          (unsigned long long int*)address;
    unsigned long long int old = *address_as_ull, assumed;
    do {
        assumed = old;
        old = atomicCAS(address_as_ull, assumed, 
                        __double_as_longlong(val + 
                        __longlong_as_double(assumed)));
    } while (assumed != old);
    return __longlong_as_double(old);
}

然后出于测试目的，我将所有图像都切换为 int，然后情况仍然相同，大部分时间都是垃圾，而一旦出现正确结果。

我需要打开一些编译标志吗？我正在使用 CMAKE 构建项目，使用

find_package(CUDA QUIET REQUIRED)

用于 CUDA 支持。以下是我设置 CUDA_NVCC_FLAGS 的方式

set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} -arch=sm_30"),

也许我错过了什么？

任何建议都将不胜感激，如果需要更多代码信息，我会更新问题。

Answer 1

所以原来解决这个问题的方法是添加下面这行来初始化dev_out指向的内存。

cudaMemcpy( dev_out, image_out, size_out * sizeof(int), cudaMemcpyHostToDevice );

我忘了初始化它，因为我以为它是一个输出变量，我在主机上初始化了它。

正如 talonmies 所说，它与 atomicAdd 完全无关。 atomicAdd 的 int 版本和 double 版本都可以完美运行。请记住在设备上初始化您的变量。