CUDA 中 if/else 块内的线程同步

Question

我想在 CUDA 中实现高斯消除。但是我在 if/else 中的线程同步有问题。

这是我的简单代码：

__device__ bool zr(float val) {
    const float zeroEpsilon = 1e-12f;
    return fabs(val) < zeroEpsilon;
}

__global__ void gauss(float* data, unsigned int size, bool* success) {
    //unsigned int len = size * (size + 1);
    extern  __shared__ float matrix[];
    __shared__ bool succ;
    __shared__ float div;
    unsigned int ridx = threadIdx.y;
    unsigned int cidx = threadIdx.x;
    unsigned int idx = (size + 1) * ridx  + cidx;
    matrix[idx] = data[idx];
    if (idx == 0)
        succ = true;
    __syncthreads();
    for (unsigned int row = 0; row < size; ++row) {
        if (ridx == row) {
            if (cidx == row) {
                div = matrix[idx];
                if (zr(div)) {
                    succ = false;
                    div = 1.0;
                }
            }
            __syncthreads();
            matrix[idx] = matrix[idx] / div;
            __syncthreads();
        }
        else {
            __syncthreads();
            __syncthreads();
        }
        if (!succ)
            break;
    }
    __syncthreads();
    if (idx == 0)
        *success = succ;
    data[idx] = matrix[idx];
    __syncthreads();
}

它是这样工作的：

1. 将矩阵复制到共享内存中。
2. 遍历行。
3. 将行除以对角线上的值。

问题出在 for 循环内部的 if/else 块内 - 死锁：

==Ocelot== PTX Emulator failed to run kernel "_Z5gaussPfjPb" with exception: 
==Ocelot== [PC 30] [thread 0] [cta 0] bar.sync 0 - barrier deadlock:
==Ocelot== context at: [PC: 59] gauss.cu:57:1 11111111111111111111
==Ocelot== context at: [PC: 50] gauss.cu:54:1 11111111111111111111
==Ocelot== context at: [PC: 33] gauss.cu:40:1 00000000000000011111
==Ocelot== context at: [PC: 30] gauss.cu:51:1 11111111111111100000

我不知道为什么会这样。当我从 if/else 块中删除同步时，它可以工作。谁能解释一下？

Answer 1

__syncthreads() 一直等待，直到一个线程块的所有线程都到达这一点并完成它们的计算。由于您的 if/else 条件，一些线程在 else-loop 中等待，一些在 if-loop 中等待，它们相互等待。但是 if 循环中的线程永远不会到达 else 循环。

Answer 2

__syncthreads() 正在这样做。

当 thread 作为指令到达 __syncthreads 时，它将阻塞/停止，当发生这种情况时，warp（32 个线程）也将阻塞，它将阻塞直到所有 threads 在同一个 block of threads已经达到了那个说法。

但是，如果同一block of threads 中的一个warp 或一个线程没有到达相同的__syncthreads 语句，它将死锁，因为至少一个thread 正在等待所有其他threads 到达相同的语句，如果没有发生，您将获得死锁。

您现在正在做的是将至少一个threads 排除在参与__syncthreads 事件之外，方法是将__syncthreads 放在 if 语句中，并非所有线程都将到达。因此，死锁。