CUDA分块矩阵乘法核无法正确处理非方阵问题求助

问题详情 #

自己写的CUDA分块矩阵乘法核，处理非方阵时结果过不了CPU串行版本的精度校验——比如跑./a.out 5 4 5就测试失败，但方阵场景（比如./a.out 4 4 4）完全正常。矩阵维度：A是m×k，B是k×n，输出C是m×n，程序输入格式为./a.out m k n。

已经查过边界条件、共享内存填充、网格块初始化这几块，但还是有部分结果没算对或者精度异常。附上完整代码和故障/正常示例输出：

完整代码 #

// 此处插入用户提供的完整CUDA代码

正常输出（4x4x4测试） #

CPU计算结果：
1.000000 2.000000 3.000000 4.000000
...（其余行省略）
GPU计算结果：
1.000000 2.000000 3.000000 4.000000
...（其余行省略）
校验通过：所有元素误差在允许范围内

故障输出（5x4x5测试） #

CPU计算结果：
0.500000 1.200000 3.100000 4.000000 5.500000
...（其余行省略）
GPU计算结果：
0.500000 0.000000 3.100000 4.000000 0.000000
...（其余行省略）
校验失败：共X个元素误差超出阈值

排查方向与修复建议 #

1. 分块循环边界处理漏了边缘块 #

非方阵场景下，分块遍历的终止条件如果只考虑整数倍分块，会导致边缘不足一个分块的元素没被计算：

• 检查核函数里遍历A、B分块的循环，把固定步长的终止改成动态取最小值：

  // 错误示例：只处理完整分块
  for (int t = 0; t < k; t += BLOCK_SIZE) {
  // 正确示例：覆盖剩余不足分块的部分
  for (int t = 0; t < k; t += BLOCK_SIZE) {
      int current_k = min(t + BLOCK_SIZE, k);
      // 基于current_k加载共享内存并计算
  }
  

• 加载共享内存时，要判断索引是否在矩阵实际范围内，避免读垃圾数据：

  if (row < m && t_col < k)
      As[row][t_col] = A[row * k + t_col];
  if (t_row < k && col < n)
      Bs[t_row][col] = B[t_row * n + col];
  

2. 全局内存索引计算出错 #

非方阵的行列索引转换很容易写错，尤其是C矩阵的地址计算：

• 核对核函数里C元素的全局索引：

  // 正确计算：C[m][n]的元素地址是 row * n + col
  int idx = row * n + col;
  // 别写成row * k + col这类错误形式
  

3. 网格/线程块维度没向上取整 #

当m或n不是BLOCK_SIZE的整数倍时，网格维度必须向上取整，保证所有元素都有线程处理：

• 检查主机端的网格初始化代码：

  // 错误示例：没向上取整，边缘元素无线程处理
  dim3 grid(m / BLOCK_SIZE, n / BLOCK_SIZE);
  // 正确示例：用整数运算向上取整
  dim3 grid((m + BLOCK_SIZE - 1) / BLOCK_SIZE, (n + BLOCK_SIZE - 1) / BLOCK_SIZE);
  

• 核函数开头要判断当前线程的row/col是否超出矩阵实际维度，超出直接返回：

  int row = blockIdx.y * blockDim.y + threadIdx.y;
  int col = blockIdx.x * blockDim.x + threadIdx.x;
  if (row >= m || col >= n) return;
  

4. 浮点精度累积差异 #

如果是误差超标而非结果完全错误，可能是CPU和GPU浮点运算顺序不同导致：

• 可以适当放宽精度校验阈值（比如从1e-6调到1e-5）
• 换成double类型计算，对比结果是否一致，判断是不是单精度累积误差的问题

验证步骤 #

1. 单独测试边缘行/列的元素，比如5x4x5里的第5行（索引4）、第5列（索引4），确认这些位置的计算逻辑
2. 在核函数里加printf，输出边缘块的共享内存数据和计算中间值，和CPU的对应步骤对比
3. 用极小非方阵测试（比如2x1x2），逐步缩小问题范围

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