問題タブ [ptx]

への変換.cuに問題があり.ptxます。私は次のように使用しnvccています：

"C:\ Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\bin\nvcc" -ptx -ccbin "C:\ Program Files (x86)\Microsoft Visual Studio 10.0\VC\bin" -o foo.ptx foo.cu

代わりに次のように表示されます。

foo.cuにあり\CUDA\v5.0\binます。

:) カーネル リソースを管理しようとしていたときに、PTX を調べることにしましたが、理解できないことがいくつかあります。これは私が書いた非常に単純なカーネルです:

次に、次を使用してコンパイルしnvcc --ptxas-options=-v -keep main.cuました。コンソールに次の出力が表示されました。

結果のptxは次のとおりです。

今、私が理解していないことがいくつかあります：

• ptx アセンブリによると、4+5+8+5=22 個のレジスタが使用されます。それでは、なぜused 2 registersコンパイル中にそれが言うのですか?
• アセンブリを見ると、threadId、blockId などのデータ型がu16. これはCUDA仕様で定義されていますか? または、これは CUDA ドライバーの異なるバージョン間で異なる可能性がありますか?
• 誰かが私にこの行を説明できますか: mul.wide.u16 %r1, %rh1, %rh2;? %r1は、代わりにが使用されるのはu32なぜですか?wideu32
• レジスターの名前はどのように選ばれますか? 私の花瓶では、その部分は理解できますが、(null) の部分%rは理解できません。データ型の長さに基づいて選択されていますか? すなわち: 16 ビットの場合、32 ビットの場合は null、64 ビットの場合は?hdhd
• カーネルの最後の 2 行をこれに置き換えるとout[idx] = in[idx];、プログラムをコンパイルすると、3 つのレジスタが使用されていると表示されます。現在、より多くのレジスタを使用するにはどうすればよいですか?

テスト カーネルが配列インデックスが範囲外かどうかをチェックしないという事実を無視してください。

どうもありがとうございました。

この質問はcuModuleLoadDataEx オプションに似ていますが、トピックを再度取り上げ、さらに情報を提供したいと思います。

cuModuleLoadDataEx を介して NV ドライバーで PTX 文字列をロードすると、すべてのオプションがまとめて無視されるように見えます。興味のある人なら誰でも直接、手間をかけずにこれを再現できるように、完全に機能する例を提供します。最初に小さな PTX カーネル (これを small.ptx として保存)、次に PTX カーネルをロードする C++ プログラム。

main.cc

ビルド (CUDA が /usr/local/cuda にインストールされていると仮定して、CUDA 5.0 を使用します):

誰かがコンパイルプロセスから賢明な情報を抽出できれば、それは素晴らしいことです! cuModuleLoadDataEx が説明されている CUDA ドライバー API のドキュメント (およびそれが受け入れるはずのオプション) http://docs.nvidia.com/cuda/cuda-driver-api/index.html

これを実行すると、ログは空になりjitTime、NV ドライバーによっても影響を受けませんでした。

編集：

JITコンパイル時間を取得することができました。ただし、ドライバーは OptVals として 32 ビット値の配列を想定しているようです。void *私のシステム64ビットにあるポインターの配列（）としてマニュアルに記載されているとおりではありません。したがって、これは機能します：

の配列で同じことを行うことはできないと思いますvoid *。次のコードは機能しません。

編集

JIT コンパイル時間を見ると、jitOptVals[0]誤解を招くものでした。コメントで述べたように、JIT コンパイラは以前の翻訳をキャッシュし、キャッシュされたコンパイルが見つかった場合、JIT コンパイル時間を更新しません。この値が変更されたかどうかを調べていたので、呼び出しがオプションをすべて無視すると仮定しました。そうではありません。それは正常に動作します。

2 つの事実: CUDA 5.0 では、後でリンクするために CUDA コードをさまざまなオブジェクト ファイルにコンパイルできます。CUDA アーキテクチャ 2.x では、関数が自動的にインライン化されなくなりました。

C/C++ ではいつものように、関数__device__ int foo()を に実装し、functions.cuそのヘッダーを に配置しましたfunctions.hu。関数fooは、他の CUDA ソース ファイルで呼び出されます。

を調べると、それがローカル メモリに流出してfunctions.ptxいることがわかります。foo()テストの目的で、私は の肉のすべてをコメントし、foo()それreturn 1; を作成しました.ptx. （関数は何もしないので、それが何であるか想像できません！）

しかし、実装をfoo()ヘッダー ファイル に移動して修飾子functions.hu を追加すると、__forceinline__ローカル メモリには何も書き込まれません。

ここで何が起こっているのですか？ CUDA がこのような単純な関数を自動的にインライン化しないのはなぜですか?

個別のヘッダーと実装ファイルの全体的なポイントは、コードの保守を容易にすることです。しかし、ヘッダーとそれらに一連の関数 (またはそれらすべて) を貼り付ける必要がある場合は__forceinline__、CUDA 5.0 のさまざまなコンパイル ユニットの目的を無効にします...

これを回避する方法はありますか？

シンプルで実際の例:

functions.cu:

上記の関数は、ローカル メモリにスピルします。

functions.ptx:

これらの 3 つの自明で最小限のカーネルについて考えてみましょう。レジスターの使用率は、私が予想するよりもはるかに高くなっています。なんで？

A:

対応するptx:

B:

対応するptx:

子：

対応するptx:

質問：

空のカーネル A と B が 2 つのレジスタを使用するのはなぜですか? CUDA は常に 1 つの暗黙的なレジスタを使用しますが、2 つの追加の明示的なレジスタが使用されるのはなぜですか?

カーネル C はさらにイライラします。10台登録？しかし、ポインタは2つしかありません。これにより、ポインター用に 2*2 = 4 個のレジスターが得られます。さらに 2 つの謎のレジスター (カーネル A とカーネル B によって提案された) がある場合でも、合計で 6 つになります。 まだ10 をはるかに下回っています。

興味がある場合は、ptxカーネル A のptxコードを次に示します。カーネル B のコードは、整数値と変数名を法として、まったく同じです。

そしてカーネルCの場合...

1. .local最初にローカルメモリ変数 ( )を宣言するのはなぜですか?
2. 2 つのポインター (関数の引数として指定) がレジスターに格納されるのはなぜですか? それらのための特別なパラメータスペースはありませんか？
3. おそらく、2 つの関数引数ポインターはレジスターに属しています。これが 2 つの.reg .b64 行を説明しています。しかし、.reg .s64ラインは何ですか？なぜそこにあるのですか？

それはさらに悪化します：

D:

与える

では、引数 (ポインター) を操作すると、レジスターが 10 から 6 に減少するのでしょうか?

PTX の学習を開始したいのですが、どこから始めればよいですか? これを行うための良い本/リソースはありますか?

これが役立つ場合は、x86/x64 ASM (多かれ少なかれ) を既に知っています。

PTX での mov 命令の仕組みがわかりません。

これは、a がレジスターまたは即値の場合、a を d に移動します。ちなみに、a がグローバル、ローカル、または共有状態空間の変数である場合、これは a のアドレスを d に移動できます。

a がグローバル メモリ内の変数であり、値が 0x1 の u64 を指しているとします... a のアドレスしか取得できないため、0x1 を d に格納するにはどうすればよいでしょうか??

アドレスの代わりに値を取得する方法がわかりません.. intel ASMのmov eax、アドレスのebx、値のmov eax、[ebx]のようなもの（ebxの逆参照）