CUDA实现奇偶数值并行归约求和遇问题求助

问题描述 #

我是CUDA编程新手，尝试实现奇偶数值分离的并行归约求和，但结果不符合预期。测试数组[5, 8, 0, -6, 2]，预期偶数和为4、奇数和为5，实际得到[8, 5]，怀疑是__shared__内存使用有误，相关代码如下：

__global__ void sumEvenOdd(int *a, int *b, int N){
    int column = blockIdx.x * blockIdx.x + threadIdx.x;

    __shared__ int s_data[2];

    if (column < N){ 
        if (a[column] % 2 == 0){
            s_data[0] += a[column];
        }
        else{
            s_data[1] += a[column];
        }
        __syncthreads();
        b[0] = s_data[0];
        b[1] = s_data[1];
    }
}

void initArray(int *a, int N){
    for (unsigned int i = 0; i < N; i++){
        a[i] = rand() % 100;
    }
}

void verify_result(int *a, int *b, int N){
    int *verify_b;
    verify_b = (int*)malloc(2 * sizeof(int));
    verify_b[0] = 0;
    verify_b[1] = 0;
    for (unsigned int i = 0; i < N; i++){
        if (a[i] % 2 == 0){
            verify_b[0] += a[i]; 
        }
        else{
            verify_b[1] += a[i];
        }
    }
    for (unsigned int i = 0; i < 2; i++){
        assert(verify_b[i] == b[i]);
    }
}

void printResult(int *a, int *b, int N){
    printf("\n");
    for (unsigned int i = 0; i < N; i++){
        printf("%d, ", a[i]);
    }
    printf("\n");
    for (unsigned int i = 0; i < 2; i++){
        printf("%d, ", b[i]);
    }
}

int main(){
 
    //Array sizes;
    int N = 5;
        
    //Size (in bytes) of matrix
    size_t bytes = N * sizeof(int);

    //Host pointers
    int *a, *b;
    
    // Allocate host memory
    a = (int*)malloc(bytes);
    b = (int*)malloc(2 * sizeof(int));

    // Initialize array
    initArray(a, N);

    // Device pointers
    int *d_a, *d_b;

    // Allocated device memory
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, 2 * sizeof(int));

    // Copy data to the device
    cudaMemcpy(d_a, a, bytes, cudaMemcpyHostToDevice);

    //Number of threads
    int THREADS = 128;

    //Number of blocks
    int BLOCKS = (N + THREADS - 1) / THREADS;

    // Launch kernel
    sumEvenOdd<<<BLOCKS, THREADS>>>(d_a, d_b, N);
    cudaDeviceSynchronize();

    // Copy back to the host
    cudaMemcpy(b, d_b, 2 * sizeof(int), cudaMemcpyDeviceToHost);

    // Check result
    verify_result(a, b, N);

    printResult(a, b, N);

    return 0;
}

错误分析 #

1. 线程索引计算错误：blockIdx.x * blockIdx.x是错误写法，正确的全局线程索引应为blockIdx.x * blockDim.x + threadIdx.x，原代码会导致线程访问错误的数组元素甚至越界。
2. 共享内存未初始化：__shared__ int s_data[2]的初始值是随机垃圾值，直接累加会导致结果失真。
3. 多线程写入竞争：多个线程同时对s_data[0]或s_data[1]执行+=操作，属于非原子操作，会引发数据竞争，结果不可预测。
4. 全局内存重复写入：每个符合条件的线程都会执行b[0] = s_data[0]和b[1] = s_data[1]，最终只有最后一个线程的写入生效，覆盖了之前的累加结果。
5. 同步点位置错误：__syncthreads()放在单个线程写入共享内存之后，此时其他线程可能还未完成写入，同步毫无意义且会打乱执行顺序。

修正后的代码 #

核心Kernel修正 #

__global__ void sumEvenOdd(int *a, int *b, int N){
    int column = blockIdx.x * blockDim.x + threadIdx.x;

    __shared__ int s_data[2];
    // 仅由块内线程0初始化共享内存为0
    if (threadIdx.x == 0) {
        s_data[0] = 0;
        s_data[1] = 0;
    }
    __syncthreads(); // 等待所有线程看到初始化后的共享内存

    if (column < N){ 
        int val = a[column];
        // 使用原子操作避免多线程写入共享内存的竞争
        if (val % 2 == 0){
            atomicAdd(&s_data[0], val);
        }
        else{
            atomicAdd(&s_data[1], val);
        }
    }
    __syncthreads(); // 等待块内所有线程完成共享内存累加

    // 仅由块内线程0将块内结果原子累加到全局内存
    if (threadIdx.x == 0) {
        atomicAdd(&b[0], s_data[0]);
        atomicAdd(&b[1], s_data[1]);
    }
}

Main函数补充修正 #

在cudaMemcpy(d_a, a, bytes, cudaMemcpyHostToDevice);之后添加设备内存初始化：

// 初始化设备端结果数组为0，确保原子累加从0开始
cudaMemset(d_b, 0, 2 * sizeof(int));

修正说明 #

• 线程索引：改用blockIdx.x * blockDim.x + threadIdx.x正确计算全局线程ID，确保每个线程访问对应的数组元素。
• 共享内存初始化：通过块内线程0初始化共享内存为0，避免随机值干扰累加结果。
• 原子操作：使用atomicAdd保证多线程对共享内存和全局内存的写入操作是原子性的，消除数据竞争。
• 同步点调整：在共享内存初始化后和块内所有线程完成累加后分别同步，确保操作顺序正确。
• 全局内存写入控制：仅让每个块的线程0执行全局内存写入，同时用原子操作累加多个块的结果，避免覆盖。
• 设备内存初始化：用cudaMemset将d_b初始化为0，确保原子累加的初始值正确。

内容的提问来源于stack exchange，提问作者Tendoskria