問題タブ [ptx]

Windows で Cuda [4.2] を使用してテクスチャを読み込むのに問題があります。

私のプログラムは、すべてのカーネル モジュールを含む ptx ファイルを読み取ります。さらに、コンパイル プロセスは、ホスト コードの短いルーチンから追加の ptx ファイルを吐き出します。ホストのみのコードを含む .cuh ファイルと .cu ファイルを次に示します。

カーネル コードで tex2d を呼び出すたびに値 0 が返されますが、確認したところ配列に有効なデータがあります。

myBind.ptx もロードする必要があるかどうか疑問に思っています。そうであれば、どのように、いつ、どこでロードされますか? またはこれは必要ありませんか？

ご協力いただきありがとうございます。

やや大きなCUDAアプリケーションがあり、達成されたGFLOPを計算する必要があります。浮動小数点演算の数を数える簡単でおそらく一般的な方法を探しています。

アセンブリ言語で事前定義されたfpoのリストを使用して、生成されたPTXコード（以下に示す）から浮動小数点演算をカウントすることは可能ですか？コードに基づいて、カウントを一般的にすることができますか？たとえば、add.s32 %r58, %r8, -2;1つの浮動小数点演算としてカウントされますか？

例：

または、FPOを数えるはるかに簡単な方法があり、これは時間の無駄ですか？

共有メモリ配列を次のように宣言する簡単なカーネルを作成しました。

カーネルの起動時に、共有メモリの number_of_bytes を指定します。カーネルを (PTX に) コンパイルしているときに、「予期しないアドレス空間」というエラーが表示されます。私はsvn(進行中の3.3)からLLVMのかなり新しいバージョンを使用しています。ここで私が間違っていることはありますか？問題は extern キーワードにあるようですが、他にどのように指定すればよいでしょうか (共有メモリ)。別の LLVM ビルドを使用する必要がありますか?

構成 CUDA 5.0 、Nvidia Tesla C1060

clang / LLVMバージョン3.2を使用して、「nbody」サンプルプログラムのカーネル（nbody_kernel.cu）のPTXコードを生成しようとしています。nbody CUDAプログラムは、NvidiaのSDKで利用できます。

https://github.com/jholewinski/llvm-ptx-samplesプロジェクトを参照しています。

以下は私のコマンドです、

clang ++ -O4 -S -I / usr / local / cuda / include -emit-llvm -target nvptx64 nbody_kernel.cu -o nbody_kernel.ll

opt -O3 -loop-unroll -unroll-allow-partial nbody_kernel.ll -o nbody_kernel.ll

llc nbody_kernel.ll -o nbody_kernel.ptx

最後のコマンド（llc）を実行した後、UNREACHABLEが実行されました！次のスタックトレースでエラーが発生しました

たとえば、threadIdx.xのようなグローバルインデックスを__builtin_ptx_read_tid_x（）などに置き換えました。LLVM IR（つまり、.ll）の生成に問題はありません。llcを使用してIRからPTXを生成しようとすると、エラーがポップアップ表示されます。

ここで何が起こっているのかについての指針はありますか？

構造体型の LLVM 値を作成しようとしています。私はLLVM-Cインターフェースを使用しており、関数を見つけています:

すべてのメンバーが LLVMConstXXX() によって作成された定数値である場合、これは正常に機能し、次のようなコードが生成されます。

しかし問題は、メンバーが定数でない場合、次のようなものが生成されることです。

そして、この LLVM コードを NVVM (Nvidia PTX バックエンド) に送信すると、次のように表示されます。

モジュール 0 (27、39): 解析エラー: 関数ローカル名の無効な使用

したがって、この構造体値の作成が正しいかどうかはわかりません。必要なのは、割り当てられたメモリではなく、値です。

誰でもアイデアがありますか？

よろしく、シャン。

CUDA ドライバー API は、ファイルシステムからの PTX コードを含むファイルの読み込みを提供します。通常、次のことを行います。

実行時に (オンザフライで) PTX ファイルを生成する場合、ファイル IO を通過するのは無駄に思えます (ドライバーがファイルを再度ロードする必要があるため)。

PTX プログラムを CUDA ドライバーに直接 (たとえば、C 文字列として) 渡す方法はありますか?

CUDA を使用した Just-In-Time (JIT) コンパイルについていくつか質問があります。

論文によると、式テンプレートに基づいてライブラリを実装しました

JM Cohen 著、「Processing Device Arrays with C++ Metaprogramming」、GPU Computing Gems - Jade Edition

それはかなりうまくいくようです。行列要素ごとの演算の計算時間を比較すると

D_D=A_D*B_D-sin(C_D)+3.;

意図的に開発された CUDA カーネルの結果を使用すると、次の結果が得られます (括弧内は行列のサイズ)。

時間 [ms] 手書きカーネル: 2.05 (1024x1024) 8.16 (2048x2048) 57.4 (4096*4096)

時間 [ミリ秒] ライブラリ: 2.07 (1024x1024) 8.17 (2048x2048) 57.4 (4096*4096)

ライブラリは、手書きのカーネルとほぼ同じ計算時間を必要とするようです。また、C++ 11 キーワード auto を使用して、実際に必要な場合にのみ式を評価しています。. 私の最初の質問は

1. (コードの最適化に関して) JIT がライブラリに提供するその他の利点はどれですか? JIT は、実行時コンパイルのためにさらに負荷を導入しますか?

Expression Templates に基づくライブラリを .dll ライブラリ内に配置できないことが知られています。 b710-825b503787e2 . 私の2番目の質問は次のとおりです。

2. JIT はサードパーティのユーザーに実装を隠すのに役立ちますか? はいの場合、どのように？

CUDA SDK には、ptx コードが実行時にロードされず、コンパイル時に定義される ptxjit の例が含まれています。私の3番目の質問は次のとおりです。

3. 私の場合、どのように JIT を実装すればよいですか? 実行時に読み込まれる PTX を使用した JIT の例はありますか?

助けてくれてありがとう。

Talonmiesのコメントに続いて編集

Cuda カーネルからのジャストインタイム (jit) コンパイルは可能ですか? 投稿、それはそれを読む

cuda コードは中間形式の ptx コードにコンパイルでき、実行時に実際のデバイス アーキテクチャ マシン コードに jit コンパイルされます。

私が持っている疑問は、上記が Expression Templates ライブラリに適用できるかどうかです。インスタンス化の問題により、CUDA/C++ テンプレート コードを PTX にコンパイルできないことはわかっています。しかし、おそらく、単項式と 2 項式の型/演算子のすべての可能な組み合わせをインスタンス化すると、実装の少なくとも一部を PTX にコンパイル (およびサード パーティのユーザーにマスク) し、次に JIT コンパイルして、建築を手元に。

CUDA で独自の 64 ビット シャッフル関数を実装しようとしています。ただし、次のようにすると:

以降の __shfl_xor へのすべての呼び出しは、引数の型に関係なく、この 64 ビット バージョンからインスタンス化されます。たとえば、私がやっている場合

それはまだ二重バージョンを使用します。回避策として、別の関数名を使用している可能性があります。しかし、このシャッフル関数をテンプレート関数から呼び出しているため、別の名前を使用すると、64 ビット浮動小数点用に別のバージョンを作成する必要があり、これはあまり適切ではありません。

では、どうすれば __shfl_xor(double,...) 関数をオーバーロードしながら、同時に __shfl_xor(int,...) を適切に呼び出すことができるのでしょうか?