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 接口的：

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

CUDA 原子添加失败

CUDA atomicAdd failed

c++

image

cuda

atomic

itk