使用SYCL子组实现矩阵乘法结果异常及GPU运行错误问题

问题分析与修复方案 #

1. 非整除维度下结果全零问题 #

当矩阵维度无法被tile_size整除时（比如M=17、tile_size=9），核心问题是边缘不完整tile的计算被遗漏：

• 错误逻辑：仅创建M/tile_size、N/tile_size个工作组，只处理能被tile_size完全覆盖的行/列，边缘剩余的行/列未被计算，导致对应位置的C矩阵保持初始全零值。
• 修复步骤：
  • 计算工作组数量时向上取整，确保覆盖所有行/列：

    size_t num_groups_row = (M + tile_size - 1) / tile_size;
    size_t num_groups_col = (N + tile_size - 1) / tile_size;
    

  • 核函数内部添加边界检查，跳过超出矩阵范围的线程：

    size_t global_i = get_global_id(0);
    size_t global_j = get_global_id(1);
    // 仅处理矩阵有效范围内的元素
    if (global_i >= M || global_j >= N) {
        return;
    }
    // 执行矩阵乘法计算逻辑
    

  • 加载tile数据时避免越界访问：
    对超出矩阵维度的位置赋值0，防止非法内存访问影响计算：

    size_t local_i = get_local_id(0);
    // 加载A矩阵的tile行
    for (size_t k = 0; k < K; ++k) {
        tile_a[local_i][k] = (global_i < M && k < K) ? A[global_i * K + k] : 0.0f;
    }
    

2. NVIDIA GPU非均匀工作组错误 #

NVIDIA CUDA设备（通过SYCL后端）不支持非均匀工作组，即所有工作组的线程数量必须完全一致：

• 错误原因：当tile_size无法整除矩阵维度时，若直接用range<2>(M, N)定义全局范围，最后一个工作组的线程数会少于其他工作组，触发non-uniform work-groups are not supported错误。
• 修复步骤：
  • 将全局范围向上取整到tile_size的倍数，确保每个工作组的大小都是tile_size × tile_size：

    range<2> global_range(num_groups_row * tile_size, num_groups_col * tile_size);
    range<2> local_range(tile_size, tile_size);
    // 使用nd_range提交核函数，明确全局与局部范围
    queue.submit([&](handler& h) {
        h.parallel_for(nd_range<2>(global_range, local_range), [=](nd_item<2> item) {
            // 核函数逻辑
        });
    });
    

  • 结合上述核函数的边界检查，让超出矩阵范围的线程直接返回，不执行计算。

验证方案 #

针对M=17、K=18、N=19、tile_size=9的场景：

1. 全局行范围设为2×9=18，全局列范围设为3×9=27
2. 核函数中global_i >=17或global_j >=19的线程直接跳过
3. 加载tile时对超出M/K/N的位置赋值0，保证计算正确性

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