問題タブ [gpu-local-memory]

ローカル メモリは、スレッドごとに 2 つのタイプであるレジスタ メモリよりも遅いというコメントしか見つかりませんでした。

共有メモリは高速であるはずですが、[スレッドの] ローカル メモリよりも高速ですか?

私がやりたいのは、メディアン フィルターのようなものですが、メディアンの代わりに特定のパーセンタイルを使用します。したがって、リストのチャンクを取得してソートし、適切なものを選択する必要があります。しかし、共有メモリ リストの並べ替えを開始できないか、問題が発生します。ローカル メモリにコピーするだけでパフォーマンスが大幅に低下しますか?

各スレッドの CUDA カーネル関数でローカル変数を宣言するとします。

また、宣言された変数がコンパイラによってローカルメモリに配置されたとします（これは、私が知る限り、1つのスレッドでのみ表示されることを除いて、グローバルメモリと同じです）。私の質問は、それを読むときにアクセスfが合体することですか?

CUDAで小さなプログラムを開発しようとしていますが、遅いのでテストをして少しググってみました。単一の変数はデフォルトでローカルスレッドメモリ内に格納されますが、配列は通常は格納されないことがわかりました。そういうわけで、実行するのにとても時間がかかるのだと思います。ここで、ローカルスレッドメモリは少なくとも16KBである必要があり、配列は52文字の長さであるため、ローカルメモリに格納する方法はありますか（構文をお願いします:)）。

次のようなものではありません：

ローカルメモリについて言及しているCUDAのドキュメントを読みました。（これは主に初期のドキュメントです。）device-propertiesは、local-memサイズ（スレッドごと）を報告します。「ローカル」メモリとはどういう意味ですか？「ローカル」メモリとは何ですか？「ローカル」メモリはどこにありますか？「ローカル」メモリにアクセスするにはどうすればよいですか？__device__記憶ですよね？

device-propertiesは、グローバル、共有、および一定のmemサイズも報告します。これらのステートメントは正しいですか： グローバルメモリは__device__メモリです。グリッドスコープとグリッド（カーネル）の存続期間があります。 コンスタントメモリは__device__ __constant__メモリです。グリッドスコープとグリッド（カーネル）の有効期間があります。 共有メモリは__device__ __shared__メモリです。単一のブロックスコープと（スレッドの）そのブロックの存続期間があります。

共有メモリはSMメモリだと思います。つまり、その単一のSMのみが直接アクセスできるメモリ。かなり限られたリソース。SMは一度にたくさんのブロックを割り当てられていませんか？これは、SMが異なるブロックの実行をインターリーブできる（またはインターリーブできない）ことを意味しますか？つまり、ブロック* A *スレッドをストールするまで実行します。次に、ブロック* B *スレッドを停止するまで実行します。次に、再びブロック* A *スレッドにスワップバックします。またはSMは、ブロック* A *のスレッドのセットを、ストールするまで実行しますか。次に、別のブロック* A *スレッドのセットがスワップインされます。このスワップは、ブロック* A *が使い果たされるまで続きます。それからそしてその時だけ、作業はブロック* Bで始まります*。共有メモリのためにお願いします。1つのSMが2つの異なるブロックからコードをスワップインしている場合、SMはどのようにして共有メモリチャンクをすばやくスワップイン/スワップアウトしますか？（後のsenerioは真であり、共有メモリスペースのスワップイン/スワップアウトはないと思います。Block* A *は完了するまで実行され、次にblock * B *が実行を開始します。注：block * A *は異なる場合があります。ブロックよりカーネル* B *。）

CUDAの共有/ローカルメモリに関しては、まだ少しわかりません。現在、私はカーネルを持っています。カーネル内で、各スレッドはリストオブジェクトを割り当てます。このようなもの

私の現在の理解では、各スレッドはdlistローカルメモリに保存されますが、これは本当ですか？その場合、カーネルの実行の最後にdlist（別のカーネルから）各オブジェクトを取得する方法はありますか、それとも__shared__最初のスレッドによって割り当てられた動的リストの配列を使用する必要がありますか？

少し複雑すぎるかもしれませんが、リストを変更する必要はありません。私が達成しようとしている実行は、次のようになります。

1. リストを作成する（GPUでのみ実行）
2. 各リストから出力を生成します（GPUで、各スレッドによって実行され、そのスレッドに割り当てられたリストからの情報のみが必要です）。
3. リストの変更/交換（GPUで引き続き実行）
4. ホストでいくつかのブレーク条件が満たされるまで、2と3を繰り返します

前もって感謝します！

多くのレジスタを使用し、それらをローカルメモリに大量に流出させるカーネルがあります。

こぼれが非常に多いように見えるので、L1 または L2 キャッシュさえも通過すると思います。ローカル メモリは各スレッド専用なので、ローカル メモリへのアクセスはコンパイラによってどのように結合されるのでしょうか? このメモリは、グローバル メモリのように 128 バイトのトランザクションで読み取られますか? この量の流出により、メモリ帯域幅の使用率が低くなります (50%)。ピークメモリ帯域幅の最大 80% を取得する流出のない同様のカーネルがあります。

EDITnvprofツールを使用して、さらにいくつかのメトリックを抽出しました。ここで説明した手法をよく理解している場合、レジスタのスピル (4 * l1 ヒットとミス / L2 = の 4 セクタにわたるすべての書き込みの合計) により、かなりの量のメモリ トラフィックが発生します(4 * (45936 + 4278911)) / (5425005 + 5430832 + 5442361 + 5429185) = 79.6%。私がここにいるかどうか誰か確認してくれませんか？

編集

私が考えていたビット トランザクションではなく、128 バイトであることに気付きました。

CUDA で簡単な関数を作成しました。画像のサイズを 2 倍にリサイズします。1920*1080 の画像の場合、この関数は完了するまでに最大 20 ミリ秒かかります。その機能を最適化するために、いくつかの異なる方法を試しました。そして、ローカルメモリが主な理由である可能性があることがわかりました。

画像を取得するために3つの異なる方法を試しました。

• OpenCV の Gpu モジュール
• OpenCV での GpuMat へのテクスチャ バインド
• グローバル メモリから GpuMat を直接フェッチする

それらのどれも私を少し改善することはできませんでした。

次に、nvvp を使用して理由を調べます。また、ローカル メモリのオーバーヘッドは、上記の 3 つの条件すべてで最大 95% です。

そこで、コードに目を向けて、nvcc がメモリをどのように使用しているかを調べます。次に、次のような単純な関数を見つけました。

80 バイトのスタック フレームが必要です (ローカル メモリにあります)。

そして、このような別の機能:

88 バイトのスタック フレームも必要です。

問題は、なぜ私の関数がこの単純なタスクで大量のローカル メモリとレジスタを使用するのかということです。また、OpenCV の関数がローカル メモリを使用せずに同じ関数を実行できるのはなぜですか (これは nvvp によるテストで、ローカル メモリの負荷はゼロです)。

私のコードはデバッグモードでコンパイルされています。そして私のカードはGT650(192 SP/SM, 2 SM)です。

スレッドごとのローカル メモリの量が私のGPU（GTX 580、計算能力2.0）では512 Koです。

Linux で CUDA 6.5 を使用してこの制限を確認しようとしましたが失敗しました。

これが私が使用したコードです（その唯一の目的はローカルメモリの制限をテストすることであり、有用な計算は行いません）：

そしてコンパイルコマンドライン:

コンパイルは問題なく、レポートは次のとおりです。

スレッドあたり 65000 バイトに達すると、GTX 580 で実行時に「メモリ不足」エラーが発生しました。コンソールでのプログラムの正確な出力は次のとおりです。

GTX 770 GPU (CUDA 6.5 を搭載した Linux) でテストも行いました。MEMSIZE=200000 ではエラーなく実行されましたが、MEMSIZE=250000 では実行時に「メモリ不足エラー」が発生しました。

この動作を説明するにはどうすればよいですか? 私は何か間違っていますか？