CUDA：将`+=`并行应用于位于设备上的数组元素是否安全？

Question

我在编译和执行 CUDA 脚本后注意到奇怪（不正确）的行为，并且能够将其隔离为以下最小示例。首先，我为整数数组定义了一个导出到 CSV 的函数（只是为了方便调试）：

#include <stdio.h>
#include <stdlib.h>    
void int1DExportCSV(int *ptr, int n){
    FILE *f;
    f = fopen("1D IntOutput.CSV", "w");
    int i = 0;
    for (i = 0; i < n-1; i++){
        fprintf(f, "%i,", ptr[i]);
    }
    fprintf(f, "%i", ptr[n-1]);
}

然后我定义了一个核函数，它将输入数组的某个元素加一：

__global__ void kernel(int *ptr){
    int x = blockIdx.x;
    int y = blockIdx.y;
    int offset = x + gridDim.x * y;
    ptr[offset] += 1;
}

主循环分配一个名为a的向量，分配一个空数组b，并分配一个名为a的设备副本dev_a：

#define DIM 64
int main(void){
    int *a;
    a = (int*)malloc(DIM*DIM*sizeof(int));
    int i;
    for(i = 0; i < DIM*DIM; i++){
        a[i] = 0;
    }
    int *b;
    b = (int*)malloc(DIM*DIM*sizeof(int));
    int *dev_a;
    cudaMalloc( (void**)&dev_a, sizeof(int)*DIM*DIM );
    cudaMemcpy( dev_a, a, DIM*DIM*sizeof(int), cudaMemcpyHostToDevice );

然后我将dev_a 输入到 DIM-by-DIM-by-DIM 块网格中，每个块都有 DIM 线程，将结果复制回来，然后将它们导出到 CSV：

    dim3 blocks(DIM,DIM,DIM);
    kernel<<<blocks,DIM>>>(dev_a);
    cudaMemcpy( b, dev_a, sizeof(int)*DIM*DIM, cudaMemcpyDeviceToHost );
    cudaFree(dev_a);
    int1DExportCSV(b, DIM*DIM);
}

生成的 CSV 文件的长度为 DIM*DIM，并用 DIM 填充。然而，虽然长度是正确的，但它应该用 DIM*DIM 填充，因为我本质上是在启动一个 DIM*DIM*DIM*DIM 线程超立方体，其中最后两个维度都专门用于增加设备数组dev_a 加一。

我的第一反应是怀疑 ptr[offset] += 1 步骤可能是罪魁祸首，因为多个线程可能同时执行此步骤，因此每个线程可能正在更新 ptr 的旧副本，而没有意识到存在是一堆其他线程同时在做这件事。但是，我对“CUDA 的禁忌”知之甚少，无法判断这是否合理。

硬件问题（据我所知）不是问题；我使用的是 GTX560 Ti，因此允许启动 3 维网格块，并且每个块的线程数为 64，远低于 Fermi 架构规定的最大值 1024。

我犯了一个简单的错误吗？还是我的示例中存在细微的错误？

此外，我注意到当我将 DIM 增加到 256 时，结果数组似乎填充了 290 到 430 之间的随机整数！我对这种行为完全感到困惑。

Answer 1

不，这不安全。块中的线程相互踩踏。

每个线程块中的线程都在更新内存中的相同位置：

ptr[offset] += 1;

offset 对于块中的每个线程都是相同的：

int x = blockIdx.x;
int y = blockIdx.y;
int offset = x + gridDim.x * y;

这是一个禁忌。结果未定义。 而是使用atomics:

atomicAdd(ptr+offset, 1);

或某种parallel reduction 方法。