GPU共有メモリbank競合(shared memory bank conflicts)
6719 ワード
GPU共有メモリbank衝突(shared memory bank conflicts)時間2016-11-05 21:47:58 FindSpace原文http://www.findspace.name/easycoding/1784テーマ共有メモリIntroduction
GPU上でメモリを共有するbank conflictsについてまとめた.主にReferenceから翻訳し、授業内容を簡単に説明した.共有メモリ(Shared Memory)
shared memporyはスライス上の(Cacheレベル)なので、ローカルメモリ(local memory)やグローバルメモリ(global memory)よりもずっと速いですが、実際には、shared memoryの遅延はキャッシュされていないグローバルメモリの遅延より100倍も小さい(スレッド間にbank conflictsがない場合).同じblockのスレッドでshared memoryを共有します.スレッドは、同じblock内の他のスレッドにアクセスしてshared memoryにグローバルメモリからロードされたデータにアクセスできます.この機能(スレッド同期、thread synchronization)は、ユーザ管理のデータcacheを実現したり、高性能のパラレルコラボレーションアルゴリズム(パラレルコンプライアンス、parallel reductionなど)を実現したりするなど、多くの役割を果たしています.bankとは
bankは区分方式です.cpuでは,アクセスメモリはあるアドレスにアクセスし,アドレス上のデータを取得するが,ここではbanks数のアドレスに一度にアクセスし,これらのアドレス上のすべてのデータを取得し,異なるbankに論理的にマッピングする.メモリ読み出しの制御に似ています.共有メモリbank conflicts
メモリ高帯域幅同時アクセスを実現するためにshared memoryは、同時アクセス可能な等サイズメモリブロック(banks)に分割される.従って、メモリ読み書きnアドレスの動作は、b個の独立したbankが同時に動作するように行うことができ、これにより、有効帯域幅が1つのbankのb倍に向上する.
しかしながら、複数のスレッド要求のメモリアドレスが同じbankにマッピングされると、これらの要求はシリアル化される.ハードウェアはこれらの要求をx個の衝突のない要求シーケンスに分け,帯域幅は元のx分の1に低下する.しかし、1つのwarp内のすべてのスレッドが同じメモリアドレスにアクセスすると、1回のブロードキャスト(boardcast)が生成され、これらのリクエストは1回で完了します.コンピューティング能力2.0以上のデバイスもマルチキャスト能力を有し、同じwarp内で同じメモリアドレスにアクセスする一部のスレッドの要求に同時に応答することができる.
bank conflictsを最小化するには,メモリアドレスがbanksにどのようにマッピングされるかを理解することが重要である.shared memoryの連続する32ビットワードは連続するbanksに割り当てられ、各clock cycleの各bankの帯域幅は32 bitsである.
計算能力1.xのデバイス上のwarpsize=32、bank数は16である.1つのwarpの共有メモリ要求は2つに分けられ,1つは前半のwarp,1つは後半のwarpの要求である.
計算能力2.0のデバイス、warpsize=32、bankの数も32.これにより、メモリリクエストは前後2つに分割されなくなります.
計算能力3.xのデバイスbankのサイズはカスタマイズ可能で、cudaDeviceSetSharedMemConfig()はcudaSharedMemBankSizeFourByteの4バイトまたはcudaSharedMemBankSizeEightByteに構成されています.8バイトに設定すると、二重精度データのbank conflictsを効果的に回避できます.サンプル1
warpsizeが8、bank数が8とする.
元のコード:
sdataはshared memoryに定義された配列である.
s=1の場合、すべてのスレッドがforループ内の文を1回実行すると、スレッド4がアクセスするsdata[8]とsdata[9]がbank[0]とbank[1]にマッピングされ、自身のスレッド0がアクセスするアドレスがbank[0]とbank[1]にマッピングされ、同じwarpでのスレッドアクセスのアドレスが同じbankにマッピングされ、シリアル処理を余儀なくされ、bank conflictsが出現する.
変更後:
は、1サイクルで2回sdataにアクセスしたため、2回に分けてアクセスせざるを得なかったが、すべてのスレッドアクセスアドレスにアクセスするたびに8つのbank内にマッピングされ、競合がなく、最高帯域幅に達した.実験結果
サンプル2
warp size=32、banks=16、(計算能力1.xのデバイス)データマッピング関係は以下の通りである:
2-way bank conflictsを例に、s=2の場合、16スレッドthreadIdx.x 0-15,base=0から仮定すると,アクセス順は図のようにthread 0がshared[0],thread 1がshared[2].一方、thread 8がアクセスするデータアドレスはshared[16]であるが、indexから15のためbank 0にマッピングされ、thread 8-15とthread 0-7は同じwarp内のスレッドであるが、1つのbankが同時に1つのthreadにしか与えられないため、アクセスはシリアルになる必要がある.すなわちthread 0-7が先にアクセスし、thread 8-15がアクセスする.
GPU上でメモリを共有するbank conflictsについてまとめた.主にReferenceから翻訳し、授業内容を簡単に説明した.共有メモリ(Shared Memory)
shared memporyはスライス上の(Cacheレベル)なので、ローカルメモリ(local memory)やグローバルメモリ(global memory)よりもずっと速いですが、実際には、shared memoryの遅延はキャッシュされていないグローバルメモリの遅延より100倍も小さい(スレッド間にbank conflictsがない場合).同じblockのスレッドでshared memoryを共有します.スレッドは、同じblock内の他のスレッドにアクセスしてshared memoryにグローバルメモリからロードされたデータにアクセスできます.この機能(スレッド同期、thread synchronization)は、ユーザ管理のデータcacheを実現したり、高性能のパラレルコラボレーションアルゴリズム(パラレルコンプライアンス、parallel reductionなど)を実現したりするなど、多くの役割を果たしています.bankとは
bankは区分方式です.cpuでは,アクセスメモリはあるアドレスにアクセスし,アドレス上のデータを取得するが,ここではbanks数のアドレスに一度にアクセスし,これらのアドレス上のすべてのデータを取得し,異なるbankに論理的にマッピングする.メモリ読み出しの制御に似ています.共有メモリbank conflicts
メモリ高帯域幅同時アクセスを実現するためにshared memoryは、同時アクセス可能な等サイズメモリブロック(banks)に分割される.従って、メモリ読み書きnアドレスの動作は、b個の独立したbankが同時に動作するように行うことができ、これにより、有効帯域幅が1つのbankのb倍に向上する.
しかしながら、複数のスレッド要求のメモリアドレスが同じbankにマッピングされると、これらの要求はシリアル化される.ハードウェアはこれらの要求をx個の衝突のない要求シーケンスに分け,帯域幅は元のx分の1に低下する.しかし、1つのwarp内のすべてのスレッドが同じメモリアドレスにアクセスすると、1回のブロードキャスト(boardcast)が生成され、これらのリクエストは1回で完了します.コンピューティング能力2.0以上のデバイスもマルチキャスト能力を有し、同じwarp内で同じメモリアドレスにアクセスする一部のスレッドの要求に同時に応答することができる.
bank conflictsを最小化するには,メモリアドレスがbanksにどのようにマッピングされるかを理解することが重要である.shared memoryの連続する32ビットワードは連続するbanksに割り当てられ、各clock cycleの各bankの帯域幅は32 bitsである.
計算能力1.xのデバイス上のwarpsize=32、bank数は16である.1つのwarpの共有メモリ要求は2つに分けられ,1つは前半のwarp,1つは後半のwarpの要求である.
計算能力2.0のデバイス、warpsize=32、bankの数も32.これにより、メモリリクエストは前後2つに分割されなくなります.
計算能力3.xのデバイスbankのサイズはカスタマイズ可能で、cudaDeviceSetSharedMemConfig()はcudaSharedMemBankSizeFourByteの4バイトまたはcudaSharedMemBankSizeEightByteに構成されています.8バイトに設定すると、二重精度データのbank conflictsを効果的に回避できます.サンプル1
warpsizeが8、bank数が8とする.
元のコード:
__global__ void reduce0(int *g_idata, int *g_odata) {
extern __shared__ int sdata[];
// 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();
// do reduction in shared mem
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();
// write result for this block to global mem
if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}
sdataはshared memoryに定義された配列である.
s=1の場合、すべてのスレッドがforループ内の文を1回実行すると、スレッド4がアクセスするsdata[8]とsdata[9]がbank[0]とbank[1]にマッピングされ、自身のスレッド0がアクセスするアドレスがbank[0]とbank[1]にマッピングされ、同じwarpでのスレッドアクセスのアドレスが同じbankにマッピングされ、シリアル処理を余儀なくされ、bank conflictsが出現する.
変更後:
for (unsigned int s = blockDim.x/2; s > 0; s >>= 1){
if (tid < s){
sdata[tid] += sdata[tid + s];
}
__syncthreads();
}
は、1サイクルで2回sdataにアクセスしたため、2回に分けてアクセスせざるを得なかったが、すべてのスレッドアクセスアドレスにアクセスするたびに8つのbank内にマッピングされ、競合がなく、最高帯域幅に達した.実験結果
サンプル2warp size=32、banks=16、(計算能力1.xのデバイス)データマッピング関係は以下の通りである:
2-way bank conflictsを例に、s=2の場合、16スレッドthreadIdx.x 0-15,base=0から仮定すると,アクセス順は図のようにthread 0がshared[0],thread 1がshared[2].一方、thread 8がアクセスするデータアドレスはshared[16]であるが、indexから15のためbank 0にマッピングされ、thread 8-15とthread 0-7は同じwarp内のスレッドであるが、1つのbankが同時に1つのthreadにしか与えられないため、アクセスはシリアルになる必要がある.すなわちthread 0-7が先にアクセスし、thread 8-15がアクセスする.