問題タブ [nvprof]

CUDA プログラムのメトリック/イベントを測定するために、次のようなコマンド ラインを使用してみました。

Visual プロファイラーでも同じ指標を測定しましたnvvp。得られる値に違いはありませんでした。

のようなメトリックを選択すると、出力の違いに気付きましたachieved_occupancy。nvvpしかし、これは実行ごとに異なるため、またはを使用しているかどうかに関係なく、実行するたびに異なる結果が得られるのはおそらくそのためですnvprof。

質問：

nvvpとnvprofはまったく同じで、使いやすさのために のnvvp上に構築された GUI に過ぎないという印象を受けました。nvprofしかし、私はこのアドバイスを与えられました：

常にビジュアル プロファイラーを使用してください。コマンドラインは絶対に使用しないでください。

また、この質問は次のように述べています。

グローバルなロード/ストア効率、リプレイ、DRAM 使用率が必要なため、コマンド ライン プロファイラーは使用したくありません。これらはビジュアル プロファイラーでより明確に表示されます。

のような「動的な」指標を除けばachieved_occupancy、結果の違いに気づいたことはありません。それで、このアドバイスは有効ですか？nvprof働き方になんらかの不備があるのでしょうか。コマンド ライン フォームよりもビジュアル プロファイラーを使用する利点があれば教えてください。

より具体的には、nvprof間違った結果をもたらすメトリクスはありますか?

ノート：

これらはNsight との違いについて尋ねているため、私の質問はこれまたはこれと同じではありません。nvvp

OpenACC を備えた Fortran MPI コードがあります。ビッグコードです。ここで意味のあるスニペットを提供することはできません。Cray aprun で問題なく動作します:

nvprof でプロファイリングしたい。私は試します：

コードは再び正常に実行されますが、すべてが完了したら、プロファイリング データは取得されず、メッセージのみが表示されます。

他にエラー メッセージは表示されません。何がこの動作を引き起こすのか誰にも分かりますか? Cray MPI Fortran コンパイラでコンパイルしています。私のコンパイルフラグは

cudatoolkitモジュールがロードされます。

コマンドラインから nvprof --metrics を使用して帯域幅を測定するための正しいオプションはどれですか? ピーク FLOPS のパーセンテージを取得するために flop_dp_efficiency を使用していますが、マニュアルには帯域幅測定のオプションがたくさんあるようで、何を測定しているのかよくわかりません。たとえば、dram_read、dram_write、gld_read、gld_write はすべて同じように見えます。また、両方が同時に発生すると仮定して、帯域幅を読み取り+書き込みスループットの合計として報告する必要がありますか?

編集：

図の優れた回答に基づいて、デバイスメモリからカーネルへの帯域幅はどうなりますか? カーネルからデバイス メモリへのパスの帯域幅 (読み取り + 書き込み) を最小限に抑えることを考えています。これはおそらく L2 キャッシュへのドラムです。

FLOPS と帯域幅を測定することにより、カーネルが計算またはメモリにバインドされているかどうかを判断しようとしています。

nvidia gtx980カードで、トランザクション数を観察することにより、結合されたメモリアクセスをテストするための簡単なカーネルを作成しました。カーネルは、

次のカーネル構成でこれを実行すると

カーネル内のデータ アクセスは完全に統合されており、データ型は float (4 バイト) であるため、予想されるロード/ストア トランザクションの数は次のようになります。

ロード トランザクション サイズ = 32 バイト

トランザクションごとにロードできる float の数 = 32 バイト / 4 バイト = 8

10240 のデータをロードするのに必要なトランザクション数 = 10240/8 = 1280 トランザクション

データの書き込みにも同じ量のトランザクションが予想されます。

しかし、nvprof メトリックを観察すると、以下の結果が得られました

データのロードに必要なトランザクションの 2 倍が必要な理由がわかりません。しかし、ロード/ストアの効率に関しては、両方のメトリックが 100% を示します

ここで何が欠けていますか？

データ型のサイズが変わったときにカーネルで実行される命令の数を調べてみました

カスタムサイズのデータ​​構造を取得するために、次のように構造体を作成しました。

そして、上記のデータ型配列をある配列から別の配列にコピーするためだけにカーネルを作成しました

次に、単一のブロックで 32 スレッドのみに対して上記のカーネルを呼び出しました

を使用してプログラムをプロファイリングし、nvprofをチェックするinstructions per warpと、 の値の変化に伴ってパラメータ値が変化することがわかりましたDATABYTES。

私の質問は、この命令数の増加の理由がfloatArray構造体内の配列によるものかどうかです。aカーネルでコピーを呼び出すと、構造体内の配列の各要素が実際に展開およびコピーされfloatArray、より多くの命令が作成されるためです。

単一の命令を使用してカーネル内のカスタム構造体変数をコピーする方法はありますか?

cudaMalloc() で割り当てられたエンティティ (およびそのサイズ) が現在 GPU デバイスに存在することを確認する最も簡単な方法は何でしょうか? 関数内のメモリ リークを見つけたいのですが、関数が 1 回だけ呼び出されて終了した場合はメモリ リークはありませんが (cuda-memcheck でチェック)、複数回呼び出された場合はメモリ フットプリントがどんどん大きくなります。

Nsight Visual Profiler は私が尋ねるには複雑すぎるようで、cuda-memcheck はリークを検出しませんでした!

OpenCL の世界には、キューに入れられた、送信された、開始および終了時間などのイベントのすべてのプロファイリング情報をナノ秒単位で返す関数clGetEventProfilingInfoがあります。printfいつでもその情報にアクセスできるのでとても便利です。

たとえば、PyOpenCLを使用すると、次のようなコードを書くことができます。

これは私の仕事にとって非常に有益です。

nvprofやcompanyなどの外部プロファイリング ツールを使用する代わりに、コードでそのような情報を取得することは可能ですか?