2.x デバイスとのデバイスでのバンク競合とは何ですか? CUDA C プログラミング ガイドを理解しているように、2.x デバイスでは、2 つのスレッドが同じ共有メモリ バンク内の同じ 32 ビット ワードにアクセスしても、バンクの競合は発生しません。代わりに、単語がブロードキャストされます。2 つのスレッドが同じ 32 ビット ワードを同じ共有メモリ バンクに書き込む場合、1 つのスレッドだけが成功します。
オンチップ メモリは 64 KB (共有メモリは 48 KB、L1 は 16 KB、またはその逆) であり、32 バンクで構成されているため、各バンクは 2 KB で構成されていると想定しています。したがって、2 つのスレッドが同じ共有メモリ バンク内の 2 つの異なる 32 ビット ワードにアクセスすると、バンクの競合が発生すると思います。これは正しいです？
あなたの説明は正しいです。バンク競合を発生させる可能性のあるアクセス パターンは多数ありますが、単純で一般的な例としてストライド アクセスがあります。
__shared__ int smem[512];
int tid = threadIdx.x;
x = smem[tid * 2]; // 2-way bank conflicts
y = smem[tid * 4]; // 4-way bank conflicts
z = smem[tid * 8]; // 8-way bank conflicts
// etc.
バンク ID = インデックス % 32 なので、x、y、z アクセスのアドレスのパターンを見ると、32 スレッドの各ワープで、x の場合、2 つのスレッドが各バンクにアクセスし、y の場合、2 つのスレッドがアクセスすることがわかります。 4 つのスレッドが各バンクにアクセスし、z の場合、8 つのスレッドが各バンクにアクセスします。