3

私は奇妙な問題を経験しています、少なくとも私にはそれは奇妙に見えます、そして私は誰かがそれに光を当てることができるかもしれないことを望んでいました。高速ローカルアクセスを共有メモリに依存するCUDAカーネルがあります。私の知る限り、ハーフワープ内のすべてのスレッドが同じ共有メモリバンクにアクセスすると、値はワープ内のスレッドにブロードキャストされます。また、複数のワープから同じバンクにアクセスしても、バンクの競合は発生せず、シリアル化されるだけです。これを念頭に置いて、これをテストするために小さなカーネルを作成しました（元のカーネルで問題が発生した後）。スニペットは次のとおりです。

#define NUM_VALUES 16
#define NUM_LOOPS  1024

__global__ void shared_memory_test(float *output)
{
    // Create some shared memory
    __shared__ int dm_delays[NUM_VALUES];

    // Loop over NUM_LOOPS
    float accumulator = 0;
    for(unsigned c = 0; c < NUM_LOOPS; c++)
    {
        // Force shared memory update
        for(int d = threadIdx.x; d < NUM_VALUES; d++)
            dm_delays[d] = c * d;

        // __syncthreads();
        for(int d = 0; d < NUM_VALUES; d++)
            accumulator += dm_delays[d];
}

    // Store accumulated value to global memory
    for(unsigned d = 0; d < NUM_VALUES; d++)
        output[d] = accumulator;
}

私はこれを16のブロック次元で実行しました（ワープの半分、それほど効率的ではありませんが、テスト目的のためだけです）。すべてのスレッドが同じ共有メモリバンクをアドレス指定する必要があるため、競合が発生しないようにする必要があります。しかし、その逆は本当のようです。このテストには、VisualStudio2010でParallelNsightを使用しています。

私にとってさらに不思議なのは__syncthreads、外側のループでコールのコメントを外すと、銀行の競合の数が劇的に増えるという事実です。

アイデアを与えるためのいくつかの数値（これは、16スレッドの1つのブロックを含むグリッドの場合、単一のハーフワープ、NUM_VALUES = 16、NUM_LOOPS = 1024）：

  • なし__syncthreads：4つの銀行の競合
  • と__syncthreads ：4,096の銀行の競合

これをcompute_capability3.0に設定されたGTX670で実行しています

前もって感謝します

__syncthreads更新： dm_delaysの値は変更されないため、外側のループでNUM_LOOPS読み取りがないと、コンパイラーによって最適化されていたことが指摘されました。現在、どちらの場合も常に4,096のバンク競合が発生しますが、それでも共有メモリのブロードキャスト動作ではうまく機能しません。

4

1 に答える 1

0

の値はdm_delays変更されないため、これは、が存在しない場合に、コンパイラが共有メモリへの1024読み取りを最適化する場合である可能性があります__syncthreads。そこで、__syncthreads別のスレッドによって値が変更される可能性があると想定される可能性があるため、値を何度も読み取ります。

于 2013-03-25T15:35:51.180 に答える