12

銀行の競合がどのように発生するかを理解しようとしています。
グローバル メモリにサイズ 256 の配列があり、1 つのブロックに 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];

上記のコードは銀行の競合を引き起こしますか?

4

2 に答える 2

16

これを確認する最善の方法は、「Compute Visual Profiler」を使用してコードをプロファイリングすることです。これには CUDA ツールキットが付属しています。また、GPU Gems 3には、「39.2.3 バンク競合の回避」という優れたセクションがあります。

「同じワープ内の複数のスレッドが同じバンクにアクセスする場合、ワープのすべてのスレッドが同じ 32 ビット ワード内の同じアドレスにアクセスしない限り、バンク競合が発生します」 - まず、それぞれ 4 バイト幅の 16 のメモリ バンクがあります。基本的に、共有メモリ バンク内の同じ 4 バイトからメモリを読み取るハーフ ワープにスレッドがある場合、バンクの競合やシリアライゼーションなどが発生します。

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

まず、配列がint型( 32 ビット ワード)であると仮定します。コードはこれらの int を共有メモリに保存し、K 番目のスレッドが K 番目のメモリ バンクに保存しているハーフ ワープ全体で保存します。たとえば、最初のハーフ ワープのスレッド 0 はshared_a[0]最初のメモリ バンクに保存され、スレッド 1 は に保存されshared_a[1]、各ハーフ ワープには 16 のスレッドがあり、これらは 16 の 4 バイト バンクにマップされます。次のハーフ ワープでは、最初のスレッドがその値を最初のスレッドにある shared_a[16] に保存します。メモリバンクを再び。したがって、int、float などの 4 バイト ワードを使用する場合、最初の例ではバンク コンフリクトは発生しません。char などの 1 バイトのワードを使用すると、最初の半分のワープ スレッド 0、1、2、および 3 はすべて、それらの値を共有メモリの最初のバンクに保存し、バンクの競合が発生します。

2 番目の例:

繰り返しますが、これはすべて使用しているワードのサイズに依存しますが、例では 4 バイトのワードを使用します。前半のワープを見ると、次のようになります。

スレッド数 = 32

N = 64

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

ハーフ ワープ全体のすべてのスレッドが同時に実行されるため、共有メモリへの書き込みによってバンクの競合が発生することはありません。ただし、これを再確認する必要があります。

これがお役に立てば幸いです。大きな返信で申し訳ありません。

于 2010-12-09T09:46:04.367 に答える
3

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

このコードを NVIDIA Visual Profiler でプロファイリングすると、要素サイズが 32 未満で 4 の倍数 (4、8、12、...、28) の場合、共有メモリへの連続アクセスはバンク競合を引き起こさないことが示されます。ただし、要素サイズが 32 の場合、バンク競合が発生します。

Ljdawson による回答には、古い情報が含まれています。

... char などの 1 バイトのワードを使用する場合、最初の半分のワープ スレッド 0、1、2、および 3 はすべて、それらの値を共有メモリの最初のバンクに保存し、バンクの競合が発生します。

これは古い GPU には当てはまるかもしれませんが、cc >= 2.x の最近の GPU では、事実上、ブロードキャスト メカニズムにより、バンクの競合は発生しません (リンク)。以下の引用は、CUDA C PROGRAMMING GUIDE (v8.0.61) G3.3 からのものです。共有メモリ。

ワープの共有メモリ要求は、同じ 32 ビット ワード内の任意のアドレスにアクセスする 2 つのスレッド間でバンク競合を生成しません (2 つのアドレスが同じバンクにある場合でも)。要求元のスレッドにブロードキャストされ (単一のトランザクションで複数のワードをブロードキャストできます)、書き込みアクセスの場合、各アドレスはスレッドの 1 つだけによって書き込まれます (どのスレッドが書き込みを実行するかは未定義です)。

これは、特に、c​​har の配列が次のようにアクセスされる場合、バンクの競合がないことを意味します。たとえば、次のようになります。

   extern __shared__ char shared[];
   char data = shared[BaseIndex + tid];
于 2017-04-18T10:21:31.037 に答える