4

CUDA Compute Capability 2.0（Fermi）から、グローバルメモリアクセスは768KBL2キャッシュを介して機能します。どうやら、開発者はもはやグローバルメモリバンクを気にしません。ただし、グローバルメモリは依然として非常に遅いため、適切なアクセスパターンが重要です。ここで重要なのは、L2を可能な限り使用/再利用することです。そして私の質問は、どうやって？いくつかの詳細な情報、L2がどのように機能するか、たとえばスレッドごとに100〜200要素の配列が必要な場合にグローバルメモリを整理してアクセスする方法に感謝します。

4

1 に答える 1

9

L2 キャッシュはいくつかの点で役立ちますが、グローバル メモリへのアクセスを結合する必要がなくなるわけではありません。簡単に言えば、合体アクセスとは、特定の読み取り (または書き込み) 命令に対して、ワープ内の個々のスレッドが、グローバル メモリ内の隣接する連続した場所を読み取り (または書き込み)、できれば 128 バイト境界でグループとして整列されていることを意味します。 . これにより、使用可能なメモリ帯域幅が最も効果的に利用されます。

実際には、これを達成するのは難しくありません。例えば：

int idx=threadIdx.x + (blockDim.x * blockIdx.x);
int mylocal = global_array[idx];

global_arrayグローバルメモリでcudaMallocを使用して通常の方法で割り当てられると仮定すると、ワープ内のすべてのスレッドにわたって結合された（読み取り）アクセスが提供されます。このタイプのアクセスでは、使用可能なメモリ帯域幅が 100% 使用されます。

重要なポイントは、通常、メモリ トランザクションは 128 バイトのブロックで発生するということです。これはたまたまキャッシュ ラインのサイズです。ブロック内のバイトの 1 つでも要求すると、ブロック全体が読み取られます (通常は L2 に格納されます)。後でそのブロックから他のデータを読み取る場合、他のメモリ アクティビティによって削除されていない限り、通常は L2 から処理されます。これは、次のシーケンスを意味します。

int mylocal1 = global_array[0];
int mylocal2 = global_array[1];
int mylocal3 = global_array[31];

通常、これらはすべて単一の 128 バイト ブロックから処理されます。の最初の読み取りでmylocal1、128 バイトの読み取りがトリガーされます。2 番目の読み取りmylocal2は、通常、メモリからの別の読み取りをトリガーすることによってではなく、キャッシュされた値 (L2 または L1 内) から処理されます。ただし、アルゴリズムを適切に変更できる場合は、最初の例のように、すべてのデータを複数のスレッドから連続して読み取ることをお勧めします。これは、たとえば、構造体の配列ではなく配列の構造体を使用するなど、データの巧妙な編成の問題である可能性があります。

多くの点で、これは CPU キャッシュの動作に似ています。キャッシュ ラインの概念は、キャッシュから要求を処理する動作と同様です。

Fermi L1 と L2 は、ライトバックとライトスルーをサポートできます。L1 は SM ごとに使用でき、共有メモリを使用して 16KB L1 (および 48KB SM) または 48KB L1 (および 16KB SM) のいずれかに構成可能に分割されます。L2 はデバイス全体で統合され、768KB です。

私が提供するいくつかのアドバイスは、L2 キャッシュがずさんなメモリ アクセスを修正するだけであると想定しないことです。GPU キャッシュは、CPU 上の同等のキャッシュよりもはるかに小さいため、問題が発生しやすくなります。一般的なアドバイスは、キャッシュが存在しないかのようにコーディングすることです。通常は、キャッシュ ブロッキングのような CPU 指向の戦略よりも、合体アクセスの生成にコーディング作業を集中させてから、場合によっては共有メモリを利用する方が適切です。次に、すべての状況で完全なメモリ アクセスを行うことができない避けられないケースについては、キャッシュがその利点を提供できるようにします。

利用可能なNVIDIA ウェビナーのいくつかを見ることで、より詳細なガイダンスを得ることができます。たとえば、Global Memory Usage & Strategy ウェビナー(およびスライド ) またはCUDA Shared Memory & Cache ウェビナーは、このトピックの参考になります。CUDA C Programming GuideのDevice Memory Access セクションを読むこともできます。

于 2012-12-12T08:07:48.050 に答える