同じメモリバンクにアクセスする同じスレッドが競合を2回引き起こしますか？

Question

私はベクトル削減を行うカーネルに取り組んでいます。

コード：

それは基本的に0

私は512フロート要素のブロックと、このスキームを以下だベクトルの全ての位置を加算した位置に結果を格納します

//scratch[] is a vector located in shared memory with all 512 elements 
NUM_ELEMENTS = 512; 
for(stride=NUM_ELEMENTS/2; stride>=1; stride = stride/2) { 
    if (threadIdx.x < stride){ 
    scratch[threadIdx.x] += scratch[threadIdx.x + stride]; 
    } 
    __syncthreads(); 
} 

奇妙なことは、銀行間の競合が発生することを期待しています。私はそうではありません。最初の反復では、スレッド0は同じバンクにある位置0と位置256を合計しています。スレッド1は位置1と位置257を加算していきます。

これらの操作はすべて、同じバンクから、2つの異なる値を得るために、ワープの各スレッドを必要とし、まだ、私は一切の競合を取得しない：

私は何をしないのですか？

Answer 1

バンクの競合の計算は、リクエストごとに1メモリ単位で計算されます。共有負荷（右側）と共有ストア（左側）は、別々の命令として多くのクロックサイクルで実行されます。