6

NVidiaからの2つのドキュメントに混乱しています。「CUDAのベストプラクティス」では、共有メモリはバンクで編成されており、一般に32ビットモードでは各4バイトがバンクであると説明されています（これが私が理解した方法です）。ただし、CUDAを使用したParallel Prefix Sum（Scan）では、バンクの競合のためにスキャンアルゴリズムにパディングを追加する方法について詳しく説明します。

私にとっての問題は、提示されているこのアルゴリズムの基本的なタイプはfloatであり、そのサイズは4バイトであるということです。したがって、各フロートはバンクであり、バンクの競合はありません。

だから私の理解は正しいです-つまり、4 * Nバイトのタイプで作業している場合、定義上、銀行の競合はないので、銀行の競合について心配する必要はありませんか？いいえの場合、それをどのように理解する必要がありますか（パディングを使用する場合）？

4

2 に答える 2

12

NVIDIA CUDA ウェビナー ページのこのウェビナーに興味があるかもしれません。バンクを含む共有メモリは、このウェビナーの スライド 35 ～ 45 でも説明されています。

一般に、共有メモリ バンクの競合は、2 つの異なるスレッドが (同じカーネル命令から) 下位 4 ビット (cc2.0 以前のデバイス) または 5 ビット (cc2.0 以降のデバイス) の共​​有メモリ内の場所にアクセスしようとすると発生する可能性があります。アドレスのデバイス) は同じです。バンク競合が発生すると、共有メモリ システムは同じバンク内のロケーションへのアクセスをシリアル化するため、パフォーマンスが低下します。パディングは、一部のアクセス パターンでこれを回避しようとします。cc2.0 以降では、すべてのビットが同じ (つまり、同じ位置) の場合、バンク競合は発生しないことに注意してください。

絵的には、次のように見ることができます。

__shared__ int A[2048];
int my;
my = A[0]; // A[0] is in bank 0
my = A[1]; // A[1] is in bank 1
my = A[2]; // A[2] is in bank 2
...
my = A[31]; // A[31] is in bank 31 (cc2.0 or newer device)
my = A[32]; // A[32] is in bank 0
my = A[33]; // A[33] is in bank 1

現在、ワープ内のスレッド間で共有メモリにアクセスすると、バンクの競合が発生する可能性があります。

my = A[threadIdx.x];    // no bank conflicts or serialization - handled in one trans.
my = A[threadIdx.x*2];  // 2-way bank conflicts - will cause 2 level serialization
my = A[threadIdx.x*32]; // 32-way bank conflicts - will cause 32 level serialization

上記の 2 方向のバンク コンフリクトを詳しく見てみましょう。2を乗算threadIdx.x しているため、スレッド 0 はバンク 0 のロケーション 0 にアクセスしますが、スレッド 16 はバンク 0 にあるロケーション 32 にアクセスし、バンク競合が発生します。上記の 32 ウェイの例では、すべてのアドレスがバンク 0 に対応します。したがって、共有メモリへの 32 のトランザクションは、すべてシリアル化されるため、この要求を満たすために発生する必要があります。

質問に答えるために、たとえば、アクセスパターンが次のようになることがわかっているとします。

my = A[threadIdx.x*32];

次に、データストレージをパディングしてA[32]、ダミー/パッドの場所にする場合がA[64]ありますA[96] 。次に、次のように同じデータを取得できます。

my = A[threadIdx.x*33];

そして、銀行の競合なしでデータを取得します。

お役に立てれば。

于 2013-02-24T21:51:52.880 に答える
7

あなたの理解は間違っています。バンクの競合は、同じワープのスレッドが同じバンクに存在する異なる値にアクセスしているときに発生します。

CUDA Cプログラミングガイドから：

高帯域幅を実現するために、共有メモリは、同時にアクセスできるバンクと呼ばれる同じサイズのメモリ モジュールに分割されます。したがって、n 個の異なるメモリ バンクに分類される n 個のアドレスから作成されたメモリの読み取りまたは書き込み要求は、同時に処理でき、単一モジュールの帯域幅の n 倍の全体帯域幅が得られます。

ただし、メモリ要求の 2 つのアドレスが同じメモリ バンクにある場合は、バンクの競合が発生し、アクセスをシリアル化する必要があります。ハードウェアは、バンク競合のあるメモリ要求を必要に応じて競合のない個別の要求に分割し、個別のメモリ要求の数に等しい係数でスループットを低下させます。個別のメモリ要求の数が n の場合、最初のメモリ要求は n 方向のバンク競合を引き起こすと言われます。

パディングは、バンクの競合を避けるために使用されます。共有メモリのアクセス パターンがわかっている場合は、バンクの競合を回避するために共有メモリ配列をパディングする方法を決定できます。

たとえば__shared__ float x[32][32];、スレッド インデックス tid を持つ各スレッドがこのように x にアクセスしているとしますsomevariable = x[tid][0];。これにより、すべてのスレッドが同じバンクから異なる値にアクセスするため、32 方向のバンク競合が発生します。
競合を避けるために、最初の次元の配列にもう 1 つの要素をパディングします: __shared__ float x[32][33];. これにより、バンクの競合が完全に解消されます。これは、各行が前の行に対して 1 オフセットされたバンク ロケーションを持つためです。

于 2013-02-24T21:51:03.197 に答える