7

「各ワープには連続したスレッドIDが含まれ、最初のワープにはスレッド0が含まれているため、最初の32スレッドを最初のワープに含める必要があります。また、1つのワープ内のすべてのスレッドが、使用可能なストリーミングマルチプロセッサで同時に実行されることも知っています。

私が理解したように、そのため、ワープが1つだけ実行されている場合は、スレッドの同期は必要ありません。__syncthreads()しかし、最後から2番目のifブロックのいずれかを削除すると、以下のコードは間違った答えを生成します。私は原因を見つけようとしましたが、何もできませんでした。私は本当にあなたの助けを願っています、それであなたはこのコードの何が悪いのか教えてくれますか？なぜ私は最後だけを残し__syncthreads()て正しい答えを得ることができないのですか？

#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);
...
}

PS私はGT560Tiを使用しています

4

1 に答える 1

7

共有メモリ変数をvolatileとして宣言する必要があります。

__shared__ volatile int data [BLOCK_SIZE];

あなたが見ている問題は、Fermiアーキテクチャとコンパイラの最適化のアーティファクトです。Fermiアーキテクチャには、共有メモリを直接操作するための命令がありません（G80 / 90 / GT200シリーズに存在していました）。したがって、すべてがロードされて登録され、操作され、共有メモリに保存されます。ただし、コンパイラは、一連の操作がレジスタにステージングされ、共有メモリとの間で中間のロードやストアが行われなければ、コードを高速化できると自由に推測できます。この種の削減コードのように、共有メモリを操作する同じワープ内のスレッドの暗黙的な同期に依存している場合を除いて、これは完全に問題ありません。

共有メモリバッファを揮発性として宣言することにより、リダクションの各段階の後に共有メモリの書き込みを強制するようにコンパイラに強制し、ワープ内のスレッド間の暗黙的なデータ同期が復元されます。

この問題は、CUDAツールキットに同梱されている（またはおそらく同梱されている）Fermiのプログラミングノートで説明されています。

于 2012-12-21T17:07:00.537 に答える