多CUDA GPU并行归约实现失败：双GPU点积结果始终为0

可能的原因及对应解决办法 #

1. GPU设备上下文未正确绑定 #

OpenMP多线程环境下，每个线程需要明确绑定到对应的GPU设备。若未在每个线程中调用cudaSetDevice()，会导致多线程共用同一GPU上下文，甚至出现内存访问冲突，最终归约结果被覆盖或未正确写入。

• 解决：在OpenMP线程入口处添加设备绑定代码，并检查返回值：

  #pragma omp parallel num_threads(2)
  {
      int tid = omp_get_thread_num();
      cudaError_t err = cudaSetDevice(tid);
      if (err != cudaSuccess) {
          printf("Failed to set device %d: %s\n", tid, cudaGetErrorString(err));
          return;
      }
      // 后续GPU计算逻辑
  }
  

2. 任务拆分与核函数索引越界 #

多GPU拆分向量时，若子向量的起始索引或长度计算错误，会导致核函数访问超出分配的全局内存范围，触发未定义行为（结果为0是常见表现）。另外，归约核函数未判断当前线程处理的元素是否在子向量范围内，也会导致越界。

• 解决：
  • 严格计算每个GPU的任务范围：

    int num_gpus = 2;
    int total_len = N;
    int start = tid * (total_len / num_gpus);
    int len = (tid == num_gpus - 1) ? (total_len - start) : (total_len / num_gpus);
    

  • 核函数中添加索引合法性检查：

    __global__ void dot_reduce(float *vec_a, float *vec_b, float *partial_sum, int len) {
        int idx = blockIdx.x * blockDim.x + threadIdx.x;
        if (idx >= len) return; // 避免越界访问
        // 归约逻辑
    }
    

3. 内存拷贝与核函数未同步 #

多GPU场景下，若未等待核函数执行完成就读取结果，或未同步OpenMP线程，会导致主机端读取到未初始化的内存（默认值为0）。

• 解决：
  • 每个线程启动核函数后，调用cudaDeviceSynchronize()确保计算完成：

    dot_reduce<<<grid, block>>>(d_a, d_b, d_partial, len);
    cudaDeviceSynchronize();
    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess) {
        printf("Kernel error: %s\n", cudaGetErrorString(err));
    }
    

  • 所有GPU完成计算后，再进行主机端汇总，添加OpenMP屏障：

    #pragma omp barrier
    #pragma omp master
    {
        // 汇总各GPU的部分和
    }
    

4. 归约核函数共享内存问题 #

GPU归约依赖共享内存时，若共享内存大小配置错误（如硬编码大小不匹配线程块）、未初始化共享内存，会导致数据破坏，结果异常。

• 解决：
  • 使用动态共享内存适配不同线程块大小：

    __global__ void dot_reduce(float *vec_a, float *vec_b, float *partial_sum, int len) {
        extern __shared__ float sdata[];
        int tid = threadIdx.x;
        int idx = blockIdx.x * blockDim.x + threadIdx.x;
        sdata[tid] = (idx < len) ? vec_a[idx] * vec_b[idx] : 0.0f;
        __syncthreads();
        // 归约逻辑
    }
    

  • 启动核函数时指定共享内存大小：

    int block_size = 256;
    int grid_size = (len + block_size - 1) / block_size;
    dot_reduce<<<grid_size, block_size, block_size * sizeof(float)>>>(d_a, d_b, d_partial, len);
    

快速排查步骤 #

• 给所有CUDA API调用添加错误检查，定位是否有设备初始化、内存分配/拷贝、核函数执行错误；
• 打印每个GPU处理的子向量范围、归约后的部分和，验证单GPU子任务结果是否正确；
• 暂时禁用OpenMP，手动模拟多GPU任务拆分，确认每个GPU的归约逻辑独立运行正常。

内容的提问来源于stack exchange，提问作者Ting-Kai Hsu