cuda - 並列縮小の例

Question

共有メモリを使用するスタンフォードのこの並列削減コードを見つけました。

コードは、262144 に等しい 1<<18 の要素数の例であり、正しい結果が得られます。

特定の数の要素では正しい結果が得られ、200000 や 25000 などの他の数の要素では予想とは異なる結果が得られるのはなぜですか? 私には、必要なスレッドブロックを常に指定しているように見えます

score 2 · Accepted Answer

// launch a single block to compute the sum of the partial sums
block_sum<<<1,num_blocks,num_blocks * sizeof(float)>>>

このコードはバグを引き起こします。

numblocksが13であると仮定します。

次に、カーネルのblockDim.x / 2は6になり、

if(threadIdx.x < offset)
{
    // add a partial sum upstream to our own
    sdata[threadIdx.x] += sdata[threadIdx.x + offset]; 
}

バグの原因となる最初の12個の要素のみが追加されます。

要素数が200000または250000の場合、num_blocksは奇数になり、バグが発生します。num_blocksが偶数の場合は、正常に機能します。

score 0 · Accepted Answer

このカーネルは、カーネルのブロッキングパラメータ（グリッドおよびスレッドブロックサイズ）の影響を受けます。入力サイズをカバーするのに十分なスレッドで呼び出していますか？

次の代わりに、forループを使用してこのようなカーネルを作成する方が堅牢です。

unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

何かのようなもの：

for ( size_t i = blockIdx.x*blockDim.x + threadIdx.x;
      i < N;
      i += blockDim.x*gridDim.x ) {
    sum += in[i];
}

CUDAハンドブックのソースコードには、「ブロックにとらわれない」コードの例がたくさんあります。削減コードはここにあります：

2 に答える 2