2015-03-06

CUDA で分岐中の __syncthreads の動作

CUDA で分岐中に __syncthreads() を実行させた場合にデッドロックが起きると
インターネット上で見かけた文書に書かれていた。

そこで実際に GTX 970 で実験してみたが、デッドロックは起きなかった。

__global__
void my_kernel()
{
    __shared__ int shared[2];
    shared[0] = -1;
    shared[1] = -1;

    int val;
    if (threadIdx.x == 0)
    {
        shared[1 - threadIdx.x] = threadIdx.x;
        __syncthreads();
        val = shared[threadIdx.x];
    }
    else
    {
        shared[1 - threadIdx.x] = threadIdx.x;
        __syncthreads();
        val = shared[threadIdx.x];
    }

    printf("threadIdx.x=%d, val=%d.\n", threadIdx.x, val);
}

my_kernel<<<1, 2>>>(); でカーネル関数を呼び出す。

実行結果は次のとおり。

threadIdx.x=0, val=1.
threadIdx.x=1, val=-1.

デッドロックはしなかったが、予期しない結果になった。

また、上記の関数で分岐の片方の __syncthreads() をコメントアウトしても
デッドロックはしなかった。

分岐中での __syncthreads() の動作は未定義ということだが、
これをテストした環境では何もしないという動作のようだ。

なお、上記のコードから if 文を取り除けば、
次のように正しい結果になる。

threadIdx.x=0, val=1.
threadIdx.x=1, val=0.
Labels: ,