共有（共有）メモリは、非常に高速なアクセス（グローバルメモリよりも100倍高速）による最適化の非常に効果的な手段です。 ただし、不適切に使用すると、銀行の競合が発生し、パフォーマンスが大幅に低下します。 この記事では、これらの競合がどのように発生し、どのように回避するかについて説明します。
共有メモリの競合はどのように発生しますか？
1つのワープ（デバイスバージョン2.0の場合）またはワープの半分（デバイスバージョン1.3以前の場合）からの2つ以上のストリームが、同じメモリバンクにある異なる32ビットワードに属するバイトにアクセスすると、競合が発生します。 競合する場合、アクセスはシーケンシャルです。 バンクにアクセスするスレッドの数は、競合の程度と呼ばれます。 競合の度合いがNの場合、アクセスは競合がない場合よりもN倍遅くなります。
ブロードキャストアクセスメカニズム
バージョン1.xのデバイスでは、複数のスレッドが同じバンクに属する同じワードにアクセスする場合、この要求が単一の場合にのみ競合を回避できます。この場合、ブロードキャストアクセスメカニズムが使用されます。
バージョン2.xのデバイスでは、このような要求がいくつかあり、それらは並列に実装されます（異なるスレッドがワードの異なるバイトにアクセスできます）。
デバイスバージョン2.0の機能にアクセスする
64ビットアクセスでは、ワープの半分のいずれかからの2つ以上のストリームが同じバンクに属するアドレスにアクセスする場合にのみ、バンクの競合が発生します。
128ビットアクセスでは、通常、2度目の銀行の競合が発生します。
32を超える容量のアクセスは、32、64、および128ビットの容量の要求に分割されます。
銀行へのメモリの割り当て方法
メモリは、シーケンス内の各32ビットワードがデバイスバージョン2.0の場合は32バンク、デバイスバージョン1.3以前の場合は16バンクのいずれかに順次割り当てられるように、バンクに分散されます。 したがって、銀行番号は次の式で計算できます。
銀行番号=（バイト単位のアドレス/ 4）％32-デバイスバージョン2.0の場合
銀行番号=（バイト単位のアドレス/ 4）％16-デバイスバージョン1.xの場合
競合アクセスの例
デバイスバージョン1.xの場合
1.8および16ビットアクセス
__shared__ char shmem8[32];
char data = shmem8[threadIdx.x];
この例では、最初の4バイトは同じバンクにあるため、最初の4バイトはアクセス時に競合します
この問題は、冗長データ（パディング）を追加し、アクセススキームを変更することで解決されます。
__shared__ char shmem8[32*4];
char data = shmem8[threadIdx.x*4];
16ビットアクセスの場合：
__shared__ short shmem16[32];
short data = shmem16[threadIdx.x];
この例では、最初の2つのショートは同じバンクにあるため、最初の2つのフローはアクセス時に競合します
この問題は、8ビットアクセスと同様に解決されます。
__shared__ short shmem16[32*2];
short data = shmem16[threadIdx.x*2];
2. 32ビットアクセス
このタイプのアクセスの場合、銀行の競合はそれほど明白ではありませんが、たとえば次のアクセス方式で発生する可能性があります。
__shared__ int shmem32[64];
int data1 = shmem32[threadIdx.x*2];
int data2 = shmem32[threadIdx.x*2+1];
この場合、0番目と8番目のストリームはそれぞれ0と1のバンクから読み取られるため、2次の競合が発生します。
たとえば、次のようにこの問題を解決できます。
__shared__ int shmem32_1[32];
__shared__ int shmem32_2[32];
int data1 = shmem32_1[threadIdx.x];
int data2 = shmem32_2[threadIdx.x];
デバイスバージョン2.0の場合
ブロードキャストアクセスの機能により、これらのデバイスの8ビットおよび16ビットアクセススキームはバンクの競合を引き起こしませんが、次の場合に競合が発生する可能性があります。
__shared__ int shared[64];
int data = shared[threadIdx.x*s];
sが偶数の場合、競合が発生します。
sが奇数であるが、競合が発生しない場合。
銀行の競合追跡
NVIDIA Banck Checker
競合は、CUDA Utility Toolkitの一部であるCUT_BANK_CHECKER（配列、インデックス）マクロを使用して追跡できます。 これを行うには、このマクロを使用してメモリにアクセスし、エミュレーションモードでアプリケーションを実行する必要があります。 アプリケーションが完了すると、競合レポートが印刷されます。
たとえば、次のように：
__shared__ int shared[64];
int data = CUT_BANK_CHECKER(shared, threadIdx.x*s);
CUDAプロファイラー
プロファイラーを使用して競合を追跡することもできます。 この情報は、
warp serializeセクションに表示されます。 このカウンタは、定数または共有メモリをアドレス指定するときにアクセスをシリアル化する必要があるワープの数を示します。つまり、このカウンタはバンクの競合を示します。
おわりに
結論として、銀行の競合を解決する最も効果的な方法は、プロファイラーによるアプリケーションの発生とその後の分析を最小限に抑えるアクセススキームを開発することです（これは決して不要ではありません）。