0

我想在 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 块中删除同步时，它可以工作。有人可以解释一下吗？

4

2 回答 2

2

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

于 2013-01-15T15:41:19.193 回答
1

__syncthreads()正在这样做。

当一条thread到达__syncthreads指令时，它会阻塞/停止，当这种情况发生时warp（32个线程）也会阻塞，它会阻塞，直到threads同一条指令中的所有线程block of threads都到达该语句。

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

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

于 2013-01-15T15:41:26.577 回答