共有CUDAメモリでvolatileを使用する場合

Question

volatileキーワードをCUDAカーネルの共有メモリに使用する必要がありますか？

__shared__ float products[THREADS_PER_ACTION]; 

// some computation 
products[threadIdx.x] = localSum; 

// wait for everyone to finish their computation 
__syncthreads(); 

// then a (basic, ugly) reduction: 
if (threadIdx.x == 0) { 
    float globalSum = 0.0f; 
    for (i = 0; i < THREADS_PER_ACTION; i++) 
     globalSum += products[i]; 
} 

が、私はこのケースで揮発する productsが必要です：私は、 volatileは、値をキャッシュしないように決してコンパイラに指示しますが、私の質問は共有配列と行動についてであることを理解できますか？各配列のエントリは、最後のスレッドを除いては、スレッド0によってすべてが読み取られる単一のスレッドによってのみアクセスされます。コンパイラが配列全体をキャッシュできる可能性があるため、 volatileである必要があります。要素？

ありがとうございます！

Answer 1

共有配列をvolatileと宣言していない場合、コンパイラは共有メモリ内の場所をレジスタ（そのスコープは単一のスレッドに限定されています）に配置したり、任意のスレッドに対して。これは、特定の共有要素に1つのスレッドのみからアクセスするかどうかにかかわらず、当てはまります。したがって、共有メモリをブロックのスレッド間の通信手段として使用する場合は、volatileと宣言することをお勧めします。

明らかに、各スレッドが共有メモリの要素にしかアクセスせず、別のスレッドに関連付けられていない場合は、これは問題ではなく、コンパイラの最適化によって何も破壊されません。

スレッドごとに共有メモリの要素にアクセスするコードセクションがあり、スレッド間のアクセスがわかりやすい場所で行われている場合は、memory fence functionを使用してコンパイラを強制的に実行できますレジスタに一時的に格納されている値を削除し、共有配列に戻します。だから、__threadfence_block()が役に立つかもしれないと思うかもしれませんが、あなたのケースでは、__syncthreads()already has memory-fencing functionality built inです。だから、あなたの__syncthreads()コールはスレッドの同期を強制するだけでなく、共有メモリのレジスタにキャッシュされた値を強制的に共有メモリに追い出すのに十分です。

ところで、コードの最後の部分の削減がパフォーマンス上問題となる場合は、並列削減方法を使用して速度を上げることを検討できます。