c++ - CUDA 原子添加失败

Question

以下 CUDA 内核应该为 3D 图像执行图像切片添加，即，您沿一维折叠 3D 体积并通过逐像素添加生成一个 2D 图像。 image_in 数据指针的大小为 128 * 128 * 128，它是使用函数 GetOutputBuffer() 从 ITK::Image 获得的。在阅读了 ITK 文档之后，我认为我们可以安全地假设数据指针指向图像数据的一段连续内存，没有填充。 image_out 只是一个大小为 128 * 128 的二维图像，也是从 ITK::Image 生成的。为了完整起见，我包含了有关图像的信息，但问题更多是关于 CUDA 原子的，可能非常初级。代码首先计算线程 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 接口(interface)的:

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 版本都可以完美运行。请记住在设备上初始化变量。