在 CUDA 中提高 parallel reduction 類程式效能的一個技巧就是 unroll last warp ，這在官方給出的示例 CUDA Radix Sort (Thrust Library) ，CUDA Parallel Reduction，scan 中都有涉及，在 CUDA_sample 中提到：

The included RadixSort class can sort either keyvalue pairs (with float or unsigned integer keys) or keys only. The optimized code in this sample (and also in reduction and scan) uses a technique known as warp-synchronous programming, which relies on the fact that within a warp of threads running on a CUDA GPU, all threads execute instructions synchronously. The code uses this to avoid __syncthreads() when threads within a warp are sharing data via __shared__ memory. It is important to note that for this to work correctly without race conditions on all GPUs, the shared memory used in these warp-synchronous expressions must be declared volatile. If it is not declared volatile, then in the absence of __syncthreads(), the compiler is free to delay stores to __shared__ memory and keep the data in registers (an optimization technique), which will result in incorrect execution. So please heed the use of volatile in these samples and use it in the same way in any code you derive from them.

這個技巧簡而言之就是利用了一個 warp 中的執行緒都是同時執行（同一個）指令的特點，從而避免一個 warp 內通過 shared memory 共享記憶體的執行緒間的 __syncthreads()  同步操作開銷。

需要指出來的是，為了使程式正確執行而避免 race condition，在這些 warp-synchronous 表示式中使用的共享記憶體必須被宣告為 volatile。 如果沒有被宣告為 volatile 的話，並且缺少 __syncthreads()，編譯器就會隨意延遲 __shared__ memory 中資料的儲存並且將資料放到暫存器中，這是由於編譯優化所造成的，這會導致錯誤的結果。

#define BLOCK_SIZE 128

__global__ void reduce ( int * inData, int * outData )
{
 __shared__ int data [BLOCK_SIZE]; 
 int tid = threadIdx.x; 
 int i   = blockIdx.x * blockDim.x + threadIdx.x; 

 data [tid] = inData [i] + inData [i + blockDim.x / 2 ];
 __syncthreads ();

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

 if ( tid < 32 )
 { 
  data [tid] += data [tid + 32];
  __syncthreads (); 
  data [tid] += data [tid + 16];
  __syncthreads (); 
  data [tid] += data [tid + 8];
  __syncthreads (); 
  data [tid] += data [tid + 4];
  __syncthreads (); 
  data [tid] += data [tid + 2];
  __syncthreads (); 
  data [tid] += data [tid + 1];
  __syncthreads (); 
 }
 if ( tid == 0 )
  outData [blockIdx.x] = data [0];
}

void main()
{
...
 reduce<<<dim3(128), dim3(128)>>>(dev_data, dev_res);
...
}
 

貼主的問題是，一旦移除程式碼中的任何一個倒數第二個 if 塊中的 __syncthreads() 語句，就會得到錯誤的答案。

而本人的情況類似，同樣的資料在數次執行後會給出不同的結果，這是由於 race condition 造成的。

回答中指出必須宣告共享記憶體為 volatile。這個問題是由於 Fermi 架構和編譯器優化所造成的。

__shared__ volatile int data [BLOCK_SIZE]; 

Fermi 架構問題是缺少對共享記憶體直接操作的指令，因此所有的資料都是在暫存器裡完成操作然後再存回共享記憶體的。

於此同時如果一系列的操作都是在暫存器中完成的話，編譯器可以隨意的規約這些操作來加快程式，而不是直接的從共享記憶體中載入，儲存資料。這在除非你依靠一個 warp 內執行緒間對共享記憶體操作的隱式同步（並行規約類程式）的情形外，都完全沒有任何問題。

將共享記憶體空間宣告為 volatile ，迫使編譯器在每步規約後都去寫 shared memory，那麼 warp 內執行緒間隱式的資料同步就能保持了。

該 slide 作者 Mark Harris 在另一篇相關的帖子中也提到這一點。

上面還提到了這個問題與架構有關，個人也有這樣的經歷：

膝上型電腦上使用的 GTX 850M（http://en.wikipedia.org/wiki/Maxwell_(microarchitecture)），屬於 Maxwell 架構，在沒有宣告 volatile 也沒顯示宣告同步時也沒有產生錯誤，宣告為 volatile 後結果也沒有改變，不知是編譯器沒有優化還是架構變得厲害了，能夠對共享記憶體直接操作了（Maxwell provides native shared memory atomic operations for 32-bit integers and native shared memory 32-bit and 64-bit compare-and-swap (CAS), which can be used to implement other atomic functions.）。

Volatile: 型別修飾符，敏感變數關鍵字；

編譯器認為其他執行緒可能隨時會修改變數的值，因此每次對該變數的引用都會被編譯成一次正式的記憶體讀取指令。作為指令關鍵字，確保本條指令不會因編譯器的優化而省略，且要求每次直接讀值。

一個定義為volatile的變數是說這變數可能會被意想不到地改變，這樣，編譯器就不會去假設這個變數的值了。精確地說就是，優化器在用到這個變數時必須每次都小心地重新讀取這個變數的值，而不是使用儲存在暫存器裡的備份。下面是volatile變數的幾個例子：

1）. 並行裝置的硬體暫存器（如：狀態暫存器）

2）. 一箇中斷服務子程式中會訪問到的非自動變數（Non-automatic variables)

3）. 多執行緒應用中被幾個任務共享的變數（摘自 http://baike.baidu.com/view/608706.htm）