GPUはメモリバンクの競合を共有しました

Question

私は銀行の紛争がどのように起こるかを理解しようとしています。
グローバルメモリにサイズ256の配列があり、単一のブロックに256個のスレッドがある場合、配列を共有メモリにコピーしたい場合があります。したがって、すべてのスレッドは1つの要素をコピーします。

shared_a[threadIdx.x]=global_a[threadIdx.x]

この単純なアクションは銀行の紛争につながりますか？

アレイのサイズがスレッドの数よりも大きいため、これを使用してグローバルメモリを共有メモリにコピーしていると仮定します。

tid = threadIdx.x;
for(int i=0;tid+i<N;i+=blockDim.x)
     shared_a[tid+i]=global_a[tid+i];

上記のコードは銀行の紛争につながりますか？

Answer 1

これをチェックする最良の方法は、「計算ビジュアルプロファイラー」を使用してコードをプロファイルすることです。これには、CUDAツールキットが付属しています。また、素晴らしいセクションがあります GPU GEMS 3 これについて - 「39.2.3銀行の紛争を回避」。

"同じワープ内の複数のスレッドが同じ銀行にアクセスすると、ワープのすべてのスレッドが同じ32ビットワード内で同じアドレスにアクセスしない限り、銀行の競合が発生します「 - 最初に16のメモリバンクがあり、それぞれ4バイト幅です。だから本質的に、あなたが持っている場合 ハーフワープの任意のスレッド 共有メモリバンクの同じ4バイトからメモリを読むと、銀行の紛争やシリアル化などがあります。

わかりました、あなたの最初の例:

まず、あなたの配列がタイプのことを言うと仮定しましょう int (32ビットワード）。コードはこれらのINTを共有メモリに保存し、KTHスレッドがKTHメモリバンクに保存しているハーフワープを越えて保存します。たとえば、前半のワープのスレッド0は shared_a[0] これは最初のメモリバンクにあり、スレッド1は shared_a[1], 、各ハーフワープには、これらのマップが16の4バイトバンクに16個のスレッドがあります。次のハーフワープでは、最初のスレッドがその値をShared_a [16]に保存します。 最初 再びメモリバンク。したがって、そのようなint、floatなどの4byteワードを使用する場合、最初の例は銀行の競合になりません。 Charなどの1バイト単語を使用する場合、前半のワープスレッド0、1、2、および3では、すべてが銀行の紛争を引き起こす共有メモリの最初の銀行に値を保存します。

2番目の例:

繰り返しますが、これはすべて使用している単語のサイズに依存しますが、例では4byteワードを使用します。だから前半のワープを見てください：

スレッド数= 32

n = 64

スレッド0：0、31、63スレッド1に書き込みます：1、32に書き込みます

ハーフワープ全体のすべてのスレッドは同時に実行されるため、共有されたメモリへの書き込みは銀行の紛争を引き起こすべきではありません。ただし、これを再確認する必要があります。

これが役立つことを願っています、大きな返信をお詫びします！

Answer 2

どちらの場合も、スレッドに連続したアドレスを持つ共有メモリにアクセスします。共有メモリの要素サイズに依存しますが、スレッドのワープによる共有メモリへの連続したアクセスは、「小さい」要素サイズの銀行競合にはなりません。

プロファイリング このコード Nvidia Visual Profilerでは、要素サイズが32より小さく、4の倍数（4、8、12、...、28）の倍数で、共有メモリへの連続したアクセスが銀行の紛争につながらないことを示しています。ただし、32の要素サイズは、銀行の紛争をもたらします。

---

ljdawsonによる回答には、いくつかの時代遅れの情報が含まれています。

... Charなどの1バイト単語を使用する場合、前半のワープスレッド0、1、2、および3では、すべてが銀行の紛争を引き起こす共有メモリの最初の銀行に値を保存します。

これは古いGPUに当てはまるかもしれませんが、CC> = 2.xを使用した最近のGPUでは、放送メカニズムのために事実上、銀行の紛争を引き起こしません。リンク）。次の引用はからです CUDA Cプログラミングガイド（v8.0.61）G3.3。共有メモリ.

ワープの共有メモリリクエストは、同じ32ビット単語内のアドレスにアクセスする2つのスレッド間に銀行の競合を生成しません（2つのアドレスが同じ銀行に該当しますが）：その場合、読み取りアクセスについては、単語は単語です。要求スレッド（複数の単語を1つのトランザクションでブロードキャストできます）にブロードキャストし、書き込みアクセスの場合、各アドレスは1つのスレッドのみで記述されます（スレッドは書き込みが未定義です）。

これは、特に、銀行の紛争がないことを意味します。

   extern __shared__ char shared[];
   char data = shared[BaseIndex + tid];