GPU共有メモリの実用的な例

Question

は、私はこのような配列を持っている：

data[16] = {10,1,8,-1,0,-2,3,5,-2,-3,2,7,0,11,0,2} 

私はG80のGPU上の共有メモリを使用して、この配列の削減を計算したいです。

NVIDIA文書に引用としてカーネルはそのようなものです：紙の

__global__ void reduce1(int *g_idata, int *g_odata) { 
extern __shared__ int sdata[]; 

unsigned int tid = threadIdx.x; 
unsigned int i = blockIdx.x*blockDim.x + threadIdx.x; 
sdata[tid] = g_idata[i]; 
__syncthreads(); 

// here the reduction : 

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

著者はこの方法でバンク競合の問題があると述べました。私は理解しようとしましたが、理由を理解できませんでしたか？私は銀行紛争と放送アクセスの定義を知っていますが、それでもこれを理解することはできません。

Bank Conflicts

Answer 1

G80プロセッサは、1.0の計算能力と、CUDA GPUの第一世代では、非常に古いCUDA可能GPUです。これらのデバイスは最近のCUDAバージョン（6.5以降）ではサポートされていないため、これらのデバイスのバンク構造を理解するのに必要な情報はオンラインドキュメントには含まれなくなりました。

G.3.3：

したがって、私はここにCUDA 6.5 CプログラミングガイドからCC 1.xのデバイスのために必要な情報を抜粋します。共有メモリ

共有メモリには、連続する32ビットワードが連続したバンクにマップするように構成された16個のバンクがあります。各バンクは、2クロック・サイクル当たり32ビットの帯域幅を有する。

ワープの共有メモリ要求は、個別に発行される ハーフワープごとに1つずつ、2つのメモリ要求に分割されます。結果として、ワープの前半に属するスレッドと同じワープの後半に に属するスレッドとの間にバンク が競合することはありません。

これらのデバイスでは、共有メモリは16バンク構造を持ち、各バンクは32ビットまたは4バイトの「幅」を持ちます。各銀行の幅は、例えば、intまたはfloatと同じです。従って、共有メモリのこの種に格納されるかもしれない最初の32 4バイトの量を想定することができ、そしてそれらの対応するバンク（アレイの名前sdataの代わりfを使用）：

extern __shared__ int f[]; 

index: f[0] f[1] f[2] f[3] ... f[15] f[16] f[17] f[18] f[19] ... f[31] 
bank: 0 1 2 3 ... 15  0  1  2  3 ... 15 

第16 int共有メモリ内の量はバンク0〜15に属し、共有メモリ内の次のもバンク0〜15に属します（int配列内にさらにデータがある場合など）。

今度は、バンク競合をトリガーするコードの行を見てみましょう：

for (unsigned int s=1; s < blockDim.x; s *= 2) { 
int index = 2 * s * tid; 
if (index < blockDim.x) { 
sdata[index] += sdata[index + s]; 
} 

のはsが1である上記のループを介して、最初のパスを考えてみましょう。つまり、各スレッドに対して、indexはthreadIdx.xのちょうど二重の値になるようindexは、2*1*tidであることを意味する。したがって、このために読ん

threadIdx.x: 0 1 2 3 4 5 6 7 8 9 10 11 ... 
index:  0 2 4 6 8 10 12 14 16 18 20 22 ... 
bank:  0 2 4 6 8 10 12 14 0 2 4 6 ... 

操作：

+= sdata[index + s] 

我々は持っている：

threadIdx.x: 0 1 2 3 4 5 6 7 8 9 10 11 ... 
index:  0 2 4 6 8 10 12 14 16 18 20 22 ... 
index + s: 1 3 5 7 9 11 13 15 17 19 21 23 ... 
bank:  1 3 5 7 9 11 13 15 1 3 5 7 ... 

したがって、最初の16スレッドでは、バンク1から読み込みたい2つのスレッド、バンク3から読み込みたいスレッド2つ、弱い2つのスレッドしたがって、このリードサイクルは、最初の16スレッドグループ全体で2ウェイバンクの競合に遭遇する。等回あたり、バンク0、2、4に、このよう

sdata[index] += 

を読み出し、その後、書き込みます：コードの同じ行の他の読み取りおよび書き込み操作は同様にバンク競合していることに注意してください16スレッドのグループ

この例を読んでいる他の人への注意：書かれているように、これはcc 1.xデバイスにのみ関連しています。 cc 2.xとそれ以降のデバイスでのバンク競合を示す方法論は似ているかもしれませんが、ワープの実行の違いと、これらの新しいデバイスが16ウェイバンクではなく32ウェイバンク構造構造。