9

Mark HarrisによるCUDAの並列削減の最適化の記事を読んだことがあり、それは非常に便利であることがわかりましたが、それでも1つまたは2つの概念を理解できないことがあります。18ページに書かれています：

//First add during load

// each thread loads one element from global to shared mem

unsigned int tid = threadIdx.x;

unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;

sdata[tid] = g_idata[i];
__syncthreads();

最適化されたコード：2回のロードと最初の削減の追加：

// perform first level of reduction,

// reading from global memory, writing to shared memory
unsigned int tid = threadIdx.x;                                    ...1

unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;          ...2

sdata[tid] = g_idata[i] + g_idata[i+blockDim.x];                   ...3

__syncthreads();                                                   ...4

2行目がわかりません。256個の要素があり、ブロックサイズとして128を選択した場合、なぜ2を掛けるのですか？ブロックサイズの決定方法を説明してください。

4

2 に答える 2

8

基本的には、下の図に示す操作を実行しています。

ここに画像の説明を入力してください

このコードは基本的に、図に示すように、スレッドの半分がグローバルメモリからの読み取りと共有メモリへの書き込みを実行することを示しています。

カーネルを実行し、いくつかの値を減らしたい場合は、上記のコードへのアクセスを、実行中のスレッド全体の半分に制限します。それぞれが512スレッドの4つのブロックがあり、上記のコードを最初の2つのブロックでのみ実行されるように制限すると、次のようになりますg_idate[4*512]。

unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;  

sdata[tid] = g_idata[i] + g_idata[i+blockDim.x];

それで：

thread 0 of block = 0  will copy the position 0 and 512,  
thread 1 of block = 0 position 1 and 513;
thread 511 of block = 0 position 511 and 1023;
thread 0 of block 1 position 1024 and 1536
thread 511 of block = 1 position 1535 and 2047

これは、各スレッドが位置にblockDim.x*2アクセスするために使用されます。したがって、次のブロックのスレッドがすでに計算された位置を計算しないことを保証するために、を乗算する必要があります。ii+blockDim.x2idg_idata

于 2012-11-29T15:32:06.510 に答える
1

最適化されたコードでは、最適化されていない実装の半分のブロックでカーネルを実行します。

最適化されていないコードでブロックのサイズを呼び出し、workこのサイズの半分をと呼びunit、これらのサイズが最適化されたコードでも同じ数値になるようにします。

最適化されていないコードでは、カーネルをそのままの数のスレッドで実行します。workつまり、ですblockDim = 2 * unit。各ブロックのコードは、の一部をg_idata共有メモリ内のサイズの配列にコピーするだけ2 * unitです。

最適化されたコードblockDim = unitでは、スレッドの1/2があり、共有メモリ内の配列は2分の1になっています。3行目では、最初の被加数は偶数の単位から、2番目は奇数の単位から来ています。このようにして、削減に必要なすべてのデータが考慮されます。

blockDim=256=work例：（単一ブロック、 ）を使用して最適化されていないカーネルを実行する場合unit=128、最適化されたコードには。の単一ブロックが含まれblockDim=128=unitます。このブロックはを取得するのでblockIdx=0、は重要では*2ありません。最初のスレッドはそうしg_idata[0] + g_idata[0 + 128]ます。

512個の要素があり、サイズ256（work=256、unit=128）の2つのブロックで最適化されていない場合、最適化されたコードには2つのブロックがありますが、サイズは128になります。2番目のブロック（blockIdx=1）の最初のスレッドはg_idata[2*128] + g_idata[2*128+128]。

于 2012-11-29T09:11:12.517 に答える