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()1つのスレッドブロックのすべてのスレッドがこのポイントに到達し、計算が完了するまで待機します。if / else-conditionのために、いくつかのスレッドはelse-loopで待機しており、いくつかはif-loopで待機しており、それらは互いに待機しています。ただし、ifループ内のスレッドがelseループに到達することはありません。

于 2013-01-15T15:41:19.193 に答える
1

__syncthreads()これをやっています。

が命令としてthread到達__syncthreadsすると、ブロック/ストールします。それが発生すると、warp（32スレッド）もブロックthreadsし、同じステートメント内のすべてblock of threadsがそのステートメントに到達するまでブロックします。

ただし、同じ1つのワープまたは1つのスレッドが同じステートメントにblock of threads到達しない場合、__syncthreads少なくとも1つthreadは他のすべてが同じステートメントに到達するのを待機しているため、デッドロックが発生します。threads発生しない場合は、デッドロックが発生します。

現在行っているのは、すべてのスレッドが到達するわけではないifステートメント内に配置することにより、イベントthreadsへの参加から少なくとも1つを除外することです。したがって、デッドロック。__syncthreads__syncthreads

于 2013-01-15T15:41:26.577 に答える