CUDA矩阵乘法程序在256<K<512时失效的原因排查

故障原因分析与修复方案 #

核心故障：线程同步死锁

当256 < K < 512时，Kpowered会被向上取整为512，每个CUDA block包含512个线程。归约循环中的__syncthreads()被放在if (k < Kpowered)分支内部，这会导致严重的同步异常：

• 第一次归约迭代时，Kpowered变为256，只有k < 256的线程会进入分支执行同步指令；
• k >= 256的线程直接跳过该分支，不会到达同步点；
• CUDA的__syncthreads()要求block内所有线程都到达同步点才能继续执行，未到达的线程会导致已进入分支的线程无限等待，最终触发死锁，程序无法正常结束。

次要问题：未初始化的共享内存

当K < Kpowered时，k >= K的线程未对local_vector[k]赋值，这些位置的内存是未定义的垃圾数据。归约过程中若合并的索引right_idx >= K，会读取垃圾值参与计算，导致矩阵C的结果错误。

修复代码示例 #

1. 修正同步逻辑与共享内存初始化

__global__
void gpu_matrix_multiply_RbyC(int M, int N, int Kpowered, int K, double *A, double *B_T, double* C) {
    int k = threadIdx.x;
    __shared__ double local_vector[512];

    int n = blockIdx.x % N;
    int m = blockIdx.x / N;

    // 初始化共享内存：k<K时赋值乘积，否则置0避免垃圾值干扰
    local_vector[k] = (k < K) ? (A[m * K + k] * B_T[n * K + k]) : 0.0;

    __syncthreads();

    // 归约循环：同步放在循环末尾，确保所有线程参与
    for (int s = Kpowered / 2; s > 0; s /= 2) {
        if (k < s) {
            local_vector[k] += local_vector[k + s];
        }
        __syncthreads(); // 所有线程必须同步，无论是否参与当前迭代
    }

    if (k == 0) {
        C[m * N + n] = local_vector[0]; // 直接赋值而非累加，避免重复调用时的错误
    }
}

2. 简化向上取整函数（可选）

原函数可通过位运算优化，逻辑更简洁高效：

int uproundToPowerOfTwo(int k) {
    assert(k > 0);
    if (k == 1) return 1;
    int n = 1;
    while (n < k) n <<= 1;
    return n;
}

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