Loading article…
並列スレッド実行( PTXまたはNVPTX [ 1 ] ) は、 Nvidiaの Compute Unified Device Architecture ( CUDA ) プログラミング環境で使用される低レベルの並列スレッド実行仮想マシンおよび命令セットアーキテクチャです。LLVM ベースのNvidia CUDA コンパイラ(NVCC) は、 OpenCL C および CUDA C/ C++で記述されたコードをPTX 命令 ( IL ) に変換し、グラフィックスドライバには、 PTX 命令を実行可能なバイナリコードに変換するコンパイラが含まれています。このバイナリコードは、 Nvidiaグラフィックス処理ユニット(GPU)の処理コアで実行できます。[ 2 ]インライン PTX アセンブリは、CUDA および OpenCL で使用できます。[ 3 ]
clangを使用したLLVM は、CUDA、OpenCL C/C++、SYCL C++、 OpenACC 、またはOpenMPディレクティブを指定して PTX を生成する機能も備えています。[ 1 ] [ 4 ] GNUコンパイラ コレクションも、指定された OpenACCまたはOpenMP ディレクティブをオフロードするために PTX を生成します。[ 5 ]
特定のGPUがサポートするPTX命令のコレクションは、その演算能力によって決まります。[ 1 ]
PTXは任意の大きさのプロセッサレジスタセットを使用します。コンパイラの出力はほぼ純粋な静的単一代入形式で、連続する行は一般的に連続するレジスタを参照します。プログラムは次の形式の宣言で始まります。
.reg .u32 %r < 335 >; // 符号なし32ビット整数型のレジスタ %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それぞれ、スレッドインデックス、ブロック次元、ブロックインデックス、グリッド次元が含まれます。[ 6 ]
ロード（ld）コマンドとストア（st）コマンドは、いくつかの異なる状態空間（メモリバンク）のいずれかを参照します。たとえばld.param、。状態空間は 8 つあります。[ 6 ] [ 7 ]
.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 は現在アクティブにメンテナンスされていません。[ 8 ]