CUDA削減の最適化

Question

Nvidia Reductionのすべての最適化を実行しようとしています。私は最初の4つの部分を実装しましたが、スライド番号22の部分＃5で立ち往生しています。

syncthreads（）がなくても、提供されたコードが動作する理由はわかりません。スレッドは、出力内の同じメモリ位置にアクセスします。

また、スライドでは、変数がvolatileに設定されていないとコードが機能しないことが示唆されています。どのように揮発性はその面で助けになるのですか？私がカーネルを呼びたくない場合、それをプログラムする最良の方法は何ですか？

私はこのコードを参考にしています。

__device__ void warpReduce(volatile int* sdata, int tid) { 
sdata[tid] += sdata[tid + 32]; 
sdata[tid] += sdata[tid + 16]; 
sdata[tid] += sdata[tid + 8]; 
sdata[tid] += sdata[tid + 4]; 
sdata[tid] += sdata[tid + 2]; 
sdata[tid] += sdata[tid + 1]; 
} 

for (unsigned int s=blockDim.x/2; s>32; s>>=1) { 
if (tid < s) 
sdata[tid] += sdata[tid + s]; 
__syncthreads(); 
} 

if (tid < 32) warpReduce(sdata, tid); 

ご協力いただきありがとうございます。さらに詳しい情報が必要な場合はコメントしてください。

Answer 1

コードは、ワープ同期プログラミングに依存しています。ワープ内で__syncthreads()を避けるのが一般的な方法でした。しかし、この動作は文書化されておらず、実際にはNVIDIAはその動作に依存するコードの作成を強く推奨していません。 Kepler tuning guideから

：

異なるスレッドがメモリを介して通信するプログラムで明示的な同期の欠如は、データ競合状態又は同期誤差を構成します。ワープシンクロナスプログラムは安全ではなく、CUDAコンパイラツールチェインで使用される最適化戦略の進化的な改良によって容易に解消されます。

例は、CUDAツールキットに付属のサンプルに含まれています。最近のバージョンを調べると、削減のこの部分は、ワープシャッフル演算機能（= 3.0）で実装され、古いデバイスには期待通りに__syncthreads()が使用されています。旧式のサンプル（CUDAツールキット6.0など）では、まだワープ同期技術で実装されていました。

まだワープ同期プログラミングについて知りたい場合は、this answerをお勧めします。