問題タブ [ptx]

cプログラミングでasmを使用する場合と同様に、コストのかかる実行を減らすために、CUDACコードでアセンブリコードを使用したいと思います。

出来ますか？

FermiGPUがL1またはL2キャッシュへのプリフェッチをサポートしていることを理解しています。しかし、CUDAリファレンスマニュアルには何も見つかりません。

Dues CUDAを使用すると、カーネルコードで特定のデータを特定のレベルのキャッシュにプリフェッチできますか？

からファイルを生成する方法と.ptxからファイルを生成する方法を知っていますが、最終的な実行可能ファイルを取得する方法がわかりません。.cu.cubin.ptx.

具体的には、にsample.cuコンパイルされたファイルがありますsample.ptx。次に、nvccを使用してにコンパイルsample.ptxしsample.cubinます。ただし、この.cubinファイルはホストコードなしで直接実行することはできません。.cubinファイルを元の.cuファイルにリンクして、最終的な実行可能ファイルを作成するにはどうすればよいですか？

ユーザーが を渡してthreads_per_blockカーネルを呼び出すコードがあるとします。次に、入力が有効かどうかを確認します (たとえば、計算能力 CC <2.0 の場合は <=512、CC >=2.0 の場合は 1024)。

nvcc -arch=sm_13CC2.0 を搭載したコンピューターにグラフィックス カードを搭載しているときに、ユーザーがパスしたときにコードをコンパイルするとどうなるのだろうかthreads_per_block == 1024? これは：

• 有効な入力 - 私が使用しているカードには CC2.0 があるため、または...
• CC1.3用にコンパイルしたので無効？

それともnvcc -arch=sm_13、少なくとも CC1.3 は必要ですが、より高い CC で実行すると、それらのより高い機能を使用できるということですか?

CUDA 4.0 がインストールされており、Compute Capability 2.0 (GTX 460 カード) を備えたデバイスがあります。

「cubin」ファイルと「ptx」ファイルの違いは何ですか?

cubin は GPU のネイティブ コードなので、これはマイクロ アーキテクチャ固有のものであり、ptx は JIT コンパイルを介して Fermi デバイス (Geforce GTX 460 など) で実行される中間言語です。ソース ファイルをコンパイルするときに.cu、ptx または cubin ターゲットを選択できます。cubin ファイルが必要な場合は、code=sm_20. しかし、ptx ファイルが必要な場合は、code=compute_20.

それが正しいか？

カーネルをさらに最適化するために、カーネルのPTXコードを読むことをお勧めしますか？

一例：自動ループ展開が機能したかどうかをPTXコードから確認できることを読みました。そうでない場合は、カーネルコードでループを手動で展開する必要があります。

• PTXコードの他のユースケースはありますか？
• PTXコードを調べますか？
• CUDAがカーネル用に生成したPTXコードを読み取る方法はどこで確認できますか？

新しいLLVMバックエンドを使用してC/C ++コードで注釈を付けたPTXアセンブラーを取得する方法を知っている人はいますか？

CUDA 4.0以前で簡単に取得できますが、CUDAツールキットをバージョン4.2にアップグレードした後、NVCCはすべてのフラグを拒否します。

NvidiaのインラインPTXアセンブリドキュメントに書かれているように、インラインアセンブリを使用するための文法は次のとおり asm("temp_string" : "constraint"(output) : "constraint"(input));
です。2つの例を次に示します。どちらの例にも、次の
asm("vadd.s32.s32.s32 %0, %1.h0, %2.h0;" : "=r"(v) : "r"(a), "r"(b));
asm("vadd.u32.u32.u32 %0.b0, %1, %2, %3;" : "=r"(v) : "r"(a), "r"(b), "r"(z));
ようなパラメータがあります。CUDAの公式文書を調べたところ、またはの意味について何の懸念も見つかりませんでした。私は、、、、、、を見てきました。私は推測するか、16ビット値を表しますが、バイト値を表します。誰かがこれらの正確な意味を知っていますか？h0b0%nh0b0h0h1b0b1b2b3h0h1bn

ロジャーダールからの助けをありがとう。PTX ISA 3.0を読んで、答えを見つけました。
「h」はハーフワードを意味します。h032ビットワードの下位ハーフワードを意味します。h132ビットワードの上位ハーフワードを意味します。「b」は整数バイトを意味します。b0、、はb1、32ビットワードの最初の8ビット、2番目の8ビット、3番目の8ビット、および最上位の8ビットb2を表します。b3

編集: この質問はオリジナルのやり直し版であるため、最初のいくつかの回答は関連性がなくなっている可能性があります。

インライン化を強制しないデバイス関数呼び出しが、デバイス関数内の同期にどのような影響を与えるかについて興味があります。問題の動作を示す簡単なテスト カーネルがあります。

カーネルはバッファを取得し、それをデバイス関数に渡します。共有バッファと、単一のスレッドを「ボス」スレッドとして識別するインジケータ変数も一緒に渡します。デバイス関数には分岐コードがあります。Boss スレッドは、最初に共有バッファーで単純な操作を実行するのに時間を費やし、次にグローバル バッファーに書き込みます。同期呼び出しの後、すべてのスレッドがグローバル バッファに書き込みます。カーネル呼び出しの後、ホストはグローバル バッファーの内容を出力します。コードは次のとおりです。

CUDA コード:

test_main.cu

test_kernel.cu

このコードは、test_main.cu の「cutilsafecall()」関数を利用するために CUDA SDK 内からコンパイルしましたが、SDK の外でコンパイルしたい場合はもちろん、これらを取り除くことができます。CUDA Driver/Toolkit バージョン 4.0、コンピューティング機能 2.0 でコンパイルし、コードは Fermi アーキテクチャの GeForce GTX 480 上で実行されました。

期待される出力は

0 1 2 3 ... blockDim.x-1

しかし、私が得る出力は

1 1 2 3 ... blockDim.x-1

これは、ボス スレッドが条件 "scratchBuffer[0] = 1;" を実行したことを示しているようです。すべてのスレッドが「scratchBuffer[threadIdx.x] = threadIdx.x;」を実行した後のステートメント ステートメントは __syncthreads() バリアで区切られていますが。

これは、同じワープ内のスレッドのバッファー位置にセンチネル値を書き込むようにボス スレッドが指示された場合でも発生します。センチネルは、適切な threadIdx.x ではなく、バッファーに存在する最終的な値です。

コードが期待される出力を生成するようにする変更の 1 つは、条件ステートメントを変更することです。

if(isBoss) {

に

if(IS_BOSS()) {

; つまり、発散制御変数をパラメータ レジスタに格納することから、マクロ関数で計算するように変更します。(ソース コード内の適切な行のコメントに注意してください。) 問題を突き止めるために私が焦点を当ててきたのは、この特定の変更です。'isBoss' 条件 (つまり、壊れたコード) と 'IS_BOSS()' 条件 (つまり、動作中のコード) を持つカーネルの逆アセンブルされた .cubin を見ると、命令の最も顕著な違いは、逆アセンブルされた壊れたコードの SSY 命令。

.cubin ファイルを "cuobjdump -sass test_kernel.cubin" で逆アセンブルして生成された逆アセンブル カーネルを次に示します。最初の 'EXIT' までがカーネルで、それ以降がデバイス関数です。唯一の違いは、デバイスの機能にあります。

分解されたオブジェクトコード:

「壊れた」コード

「働く」コード

「SSY」命令は作業コードに存在しますが、壊れたコードには存在しません。cuobjdump のマニュアルでは、「同期ポイントを設定します。潜在的に発散する可能性のある命令の前に使用されます」という命令について説明しています。これは、何らかの理由で、コンパイラーが壊れたコードの発散の可能性を認識していないのではないかと考えさせられます。

また、__noinline__ ディレクティブをコメント アウトすると、コードが期待どおりの出力を生成し、実際に「壊れた」バージョンと「動作する」バージョンによって生成されるアセンブリがまったく同じであることもわかりました。したがって、変数がコール スタックを介して渡された場合、その変数を使用して発散と後続の同期呼び出しを制御することはできないと思います。その場合、コンパイラは発散の可能性を認識していないようであり、したがって「SSY」命令を挿入しません。これが本当にCUDAの正当な制限であるかどうか、そしてもしそうなら、これがどこかに文書化されているかどうかは誰にも分かりますか?

前もって感謝します。

CUDA 2.0 デバイスで、特定の 1 つの変数に対してのみ L1 キャッシュを無効にする方法はありますか? すべてのメモリ操作に対してフラグ-Xptxas -dlcm=cgを追加して、コンパイル時に L1 キャッシュを無効にできることを知っています。nvccただし、特定のグローバル変数でのメモリ読み取りに対してのみキャッシュを無効にして、残りのすべてのメモリ読み取りが L1 キャッシュを通過するようにしたいと考えています。

私が Web で行った検索に基づいて、考えられる解決策は PTX アセンブリ コードを使用することです。