並列スレッド実行

並列スレッド実行（PTXまたはNVPTX ^{[1] ）は、} NvidiaのCompute Unified Device Architecture（CUDA）プログラミング環境で使用されている低レベルの並列 スレッド 実行 仮想マシンおよび命令セットアーキテクチャです。Nvidia CUDAコンパイラ（NVCC）は、 C++のような言語であるCUDAで記述されたコードをPTX命令（米国標準情報交換コード（ASCII ）テキストとして表されるアセンブリ言語）に変換し、グラフィックスドライバーにはPTX命令を実行可能なバイナリコードに変換するコンパイラが含まれており、 ^[2] Nvidiaグラフィックス処理ユニット（GPU）の処理コアで実行できます。GNUコンパイラコレクションには、 OpenMPオフロードのコンテキストでPTXを生成する基本的な機能もあります。^[3]インラインPTXアセンブリはCUDAで使用できます。^[4]

PTXは任意の大きさのプロセッサレジスタセットを使用します。コンパイラからの出力は、連続する行が通常連続するレジスタを参照する、ほぼ純粋な静的単一代入形式です。プログラムは、次の形式の宣言で始まります。

.reg .u32 %r < 335 >; // 符号なし 32 ビット整数型の 335 個のレジスタ %r0、%r1、...、%r334 を宣言します。              

これは 3 つの引数を取るアセンブリ言語であり、ほとんどすべての命令は、操作するデータ型 (符号と幅) を明示的にリストします。レジスタ名の前には % 文字が付き、定数はリテラルです。例:

shr .u64 %rd14 , %rd12 , 32 ; // 符号なし 64 ビット整数を %rd12 から 32 桁右にシフトし、結果は %rd14 になりますcvt .u64.u32 %rd142 , %r112 ; // 符号なし 32 ビット整数を 64 ビットに変換します        
      

述語レジスタはありますが、シェーダーモデル1.0のコンパイルされたコードは、これらを分岐コマンドと組み合わせてのみ使用します。条件分岐は

@ %p14 bra $label ; // $label に分岐               

命令setp.cc.typeは、適切なタイプの 2 つのレジスタを比較した結果を述語レジスタに設定します。また、64 ビット レジスタが 64 ビット レジスタ以下の場合は32 ビット レジスタを に設定するset命令もあります。それ以外の場合はが に設定されます。 set.le.u32.u64 %r101, %rd12, %rd28%r1010xffffffff%rd12%rd28%r1010x00000000

擬似レジスタを表す定義済み識別子がいくつかあります。特に、、%tid, %ntid, %ctaidには%nctaidそれぞれスレッドインデックス、ブロック次元、ブロックインデックス、グリッド次元が含まれます。^[5]

ロード（ld）およびストア（st）コマンドは、いくつかの異なる状態空間（メモリバンク）の1つを参照しますld.param。状態空間は8つあります。^[5]

.reg

レジスター

.sreg

特殊な読み取り専用プラットフォーム固有のレジスタ

.const

共有読み取り専用メモリ

.global

すべてのスレッドで共有されるグローバルメモリ

.local

各スレッド専用のローカルメモリ

.param

カーネルに渡されるパラメータ

.shared

ブロック内のスレッド間で共有されるメモリ

.tex

グローバルテクスチャメモリ（非推奨）

共有メモリは、PTX ファイル内のフォームの先頭の行で宣言されます。

.shared .align 8 .b8 pbatch_cache [ 15744 ]; // 8バイト境界に揃えた15,744バイトを定義する     

PTXでカーネルを書くには、CUDAドライバーAPIを介してPTXモジュールを明示的に登録する必要があり、これは通常、CUDAランタイムAPIとNvidiaのCUDAコンパイラであるnvccを使用するよりも面倒です。GPU Ocelotプロジェクトは、CUDAランタイムAPIカーネル呼び出しと一緒にPTXモジュールを登録するためのAPIを提供しましたが、GPU Ocelotは現在積極的にメンテナンスされていません。^[6]

• 標準ポータブル中間表現(SPIR)
• CUDAバイナリ（cubin） – ファットバイナリの一種

• NVIDIA Developer Zone の PTX ISA ページ