22

Compute Capability 2.0 (Fermi) がリリースされた後、共有メモリの使用例が残っているのではないかと考えていました。つまり、L1 にバックグラウンドでマジックを実行させるよりも、共有メモリを使用する方がよいのはどのような場合でしょうか?

共有メモリは、CC < 2.0 用に設計されたアルゴリズムを変更せずに効率的に実行できるようにするためのものですか?

共有メモリを介して共同作業を行うには、ブロック内のスレッドが共有メモリに書き込み、 と同期し__syncthreads()ます。単純にグローバル メモリに (L1 を介して) 書き込み、 と同期しないのはなぜ__threadfence_block()ですか? 後者のオプションは、値の 2 つの異なる場所に関連付ける必要がないため実装が容易であり、グローバルから共有メモリへの明示的なコピーがないため高速です。データは L1 にキャッシュされるため、スレッドはデータが実際にグローバル メモリに到達するまで待機する必要はありません。

共有メモリを使用すると、そこに置かれた値がブロックの期間中ずっとそこに残ることが保証されます。これは、十分に頻繁に使用されない場合に削除される L1 の値とは対照的です。アルゴリズムが実際に持っている使用パターンに基づいて L1 に管理させるよりも、このようなめったに使用されないデータを共有メモリにキャッシュした方がよい場合はありますか?

4

3 に答える 3

9

私の知る限り、GPU の L1 キャッシュは CPU のキャッシュと同じように動作します。したがって、「これは、十分に頻繁に使用されない場合に削除される L1 の値とは対照的です」というコメントは、私にはあまり意味がありません。

L1 キャッシュ上のデータは、十分に頻繁に使用されていない場合は削除されません。通常、以前はキャッシュになく、アドレスが既に使用されているメモリ領域に解決されるメモリ領域に対して要求が行われると、削除されます。NVidia で採用されている正確なキャッシュ アルゴリズムはわかりませんが、通常の n-way アソシアティブを想定すると、各メモリ エントリは、そのアドレスに基づいて、キャッシュ全体の小さなサブセットにしかキャッシュできません。

これもあなたの質問に答えるかもしれないと思います。共有メモリを使用すると、何をどこに保存するかを完全に制御できますが、キャッシュを使用すると、すべてが自動的に行われます。コンパイラと GPU は依然としてメモリ アクセスの最適化において非常に巧妙ですが、どのような入力が与えられるか、どのスレッドが何を行うか (特定のもちろん程度）

于 2012-07-01T03:11:54.827 に答える
1

複数のメモリ層を介してデータをキャッシュするには、常にキャッシュ コヒーレンシ プロトコルに従う必要があります。このようなプロトコルはいくつかあり、どれが最も適しているかは常にトレードオフになります。

いくつかの例を見ることができます：

GPU関連

一般にコンピューティング ユニット用

それは巨大なドメインであり、私は専門家ではないので、多くの詳細には入りたくありません. 私が指摘したいのは、共有メモリシステム (ここでは共有という用語多くのコンピューティング ユニット (CU) が同時にデータを必要とする、いわゆる GPU の共有メモリを指しません) 可能な限り高速にフェッチできるように、データをユニットの近くに保持しようとするメモリ プロトコルがあります。同じ SM (対称型マルチプロセッサ) 内の多くのスレッドが同じデータにアクセスする GPU の例では、スレッド 1 がグローバル メモリからバイトのチャンクを読み取り、次のサイクルでスレッド 2 がこれらのデータにアクセスしようとすると、効率的な実装は、スレッド 2 がデータが L1 キャッシュに既にあることを認識し、高速にアクセスできるようにすることです。これは、キャッシュ コヒーレンシ プロトコルが達成しようとしていることであり、すべての計算ユニットがキャッシュ L1、L2 などに存在するデータで最新の状態になるようにします。

ただし、スレッドを最新の状態に保つ、またはスレッドを一貫した状態に保つには、本質的にサイクルが失われるというコストがかかります。

CUDA では、メモリを L1 キャッシュではなく共有として定義することで、そのコヒーレンシ プロトコルからメモリを解放します。そのため、そのメモリ (それが何であれ、物理的には同じ素材) へのアクセスは直接的であり、コヒーレンシ プロトコルの機能を暗黙的に呼び出すことはありません。

これがどれくらい速くなるべきかはわかりません。私はそのようなベンチマークを実行しませんでしたが、このプロトコルにはもうお金を払わないので、アクセスはより速くなるはずです!

もちろん、NVIDIA GPU の共有メモリはバンクに分割されており、パフォーマンスの向上のために使用したい場合は、前にこれを確認する必要があります。その理由は、2 つのスレッドが同じバンクにアクセスするときに発生するバンクの競合であり、これによりアクセスのシリアル化が発生します... しかし、それは別の リンクです

于 2020-09-25T12:08:41.787 に答える