CUDA并行规约内核连续调用结果异常问题求助

问题现象 #

实现NVIDIA并行规约文档中的reduce0、reduce1内核时，单个内核运行结果正确（计算0~99的和为4950），但连续调用reduce0后再调用reduce1时，reduce1输出9900（预期4950）；移除reduce0调用后，reduce1结果恢复正常。

代码片段 #

reduce0内核 #

__global__ void reduce0(int *g_idata, int *g_odata) {    
  extern __shared__ int sdata[];                         
                                                         
  unsigned int tid = threadIdx.x;                        
  unsigned int i   = blockIdx.x*blockDim.x + threadIdx.x;
                                                         
  sdata[tid] = g_idata[i];                               
  __syncthreads();                                       
                                                         
  //reduction in shared memory                           
  for(unsigned int s=1; s<blockDim.x; s *= 2 ) {          
    if (tid % (2*s) == 0 ) {                             
      sdata[tid] += sdata[tid + s];                      
    }                                                    
    __syncthreads();                                     
  }                                                      
                                                         
  // write result for this bloc to global memory          
  if (tid == 0) g_odata[blockIdx.x] = sdata[0];

  // Try to initialize the shared memory
  //sdata[tid] = 0;
  //__syncthreads();
}                                                        

reduce1内核 #

__global__ void reduce1(int *g_idata, int *g_odata) {    
  extern __shared__ int sdata[];                         
                                                         
  unsigned int tid = threadIdx.x;                        
  unsigned int i   = blockIdx.x*blockDim.x + threadIdx.x;
                                                         
  sdata[tid] = g_idata[i];                               
  __syncthreads();                                       
                                                         
  //reduction in shared memory                           
  for(unsigned int s=1; s<blockDim.x; s *= 2 ) {          
    int index = 2*s*tid;                                 
                                                         
    if (index < blockDim.x) {                            
      sdata[index] += sdata[index+s];                    
    }                                                    
                                                         
    __syncthreads();                                     
  }                                                      
                                                         
  // write result for this bloc to global memory          
  if (tid == 0) g_odata[blockIdx.x] = sdata[0];           
}                                                        

主函数关键部分 #

int n_data = 100;
// 分配dev_arr，仅容纳100个int
cudaMalloc((void **)&dev_arr, n_data*sizeof(int));
// 配置线程块：512线程/块，共2块（100/512+1）
n_threads = 512;
n_block   = n_data/n_threads+1;
// 分配dev_arr_block，容纳n_block个int
cudaMalloc((void **)&dev_arr_block, n_block*sizeof(int));

// 连续调用reduce0和reduce1
reduce0<<<n_block, n_threads, n_threads*sizeof(int)>>>(dev_arr, dev_arr_block);
// ... 读取sum0 ...
reduce1<<<n_block, n_threads, n_threads*sizeof(int)>>>(dev_arr, dev_arr_block);
// ... 读取sum1 ...

问题根源 #

全局内存越界访问+内存地址重叠：

1. 主函数中dev_arr仅分配了100个int的空间，但内核配置了512线程/块，导致block0中tid >=100的线程计算出的i=tid >=100，访问dev_arr[i]属于越界访问全局内存。
2. CUDA的cudaMalloc会在连续的设备内存区域分配内存，dev_arr_block的起始地址恰好等于dev_arr的起始地址 + 100*sizeof(int)（即dev_arr[100]的地址）。
3. reduce0运行后，会将计算结果4950写入dev_arr_block[0]；此时dev_arr[100]的地址与dev_arr_block[0]完全重叠，所以reduce1运行时，tid=100的线程会读取到这个4950并写入共享内存。
4. reduce1的规约逻辑会将所有共享内存中的值合并到sdata[0]，因此最终总和为0~99的和（4950） + 4950 = 9900。

而单独运行reduce1时，dev_arr_block未被写入，dev_arr[100]是未初始化的垃圾值（恰好为0或未被规约逻辑合并到结果中），因此结果正确。

解决方案 #

在两个内核中，添加对有效数据范围的判断，超出范围的线程将共享内存初始化为0，避免越界读取垃圾值：

// 替换原sdata[tid] = g_idata[i]; 为以下代码
if (i < n_data)
    sdata[tid] = g_idata[i];
else
    sdata[tid] = 0;

注意：需要将n_data作为参数传入内核，或者定义为全局常量。

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