Parallel_Thread_Executionとは？わかりやすく解説

Parallel Thread Execution (PTX あるいは NVPTX^[1]) は、NVIDIAのCUDAプログラミング環境で使用されるGPU用の疑似アセンブリ言語である。nvccコンパイラは、C言語/C++の独自拡張であるCUDA C/C++で書かれたコードのうち、GPU上で実行されるデバイスコードをPTXに翻訳する。そしてグラフィックスドライバは、PTXをGPUのプロセッシングコア上で実行されるバイナリコードに翻訳するコンパイラを搭載している。中間表現であるPTXを利用することで、演算能力や機能 (compute capability: CC) および設計思想の異なるハードウェア上で共通して動作するCUDAプログラムを記述したり^[2]、ドライバが生成するバイナリコードを実行環境に応じて最適化したりすることが容易になる。

レジスタ

PTXは、任意の大きさのレジスタセットを使用する。コンパイラからの出力は、ほぼ純粋な単一代入形式である。通常、連続したレジスタを参照する連続した命令列を伴っている。プログラムは以下の形式の宣言で始まる。

.reg .u32 %r<335>;            // 符号なし32bit整数型の335個のレジスタを宣言する。

PTXは、最大で3つの引数を持つアセンブリ言語である。そして、ほとんど全ての命令は、命令が操作するデータのデータ型を（符号とデータ幅によって）明示するために命令の右にデータ型を並べて書く。レジスタ名は、先頭に%文字を付ける。そして、定数はそのまま書く。以下に例を示す。

shr.u64 %rd14, %rd12, 32;     // %rd12の符号なし64bit整数を32bit右シフトする。結果は%rd14に入る。
cvt.u64.u32 %rd142, %r112;    // 符号なし32bit整数を64bitへ変換する。

プレディケートレジスタもある（命令を実行するかどうかを決める条件を格納するレジスタ）。しかし、シェーダーモデル1.0でコンパイルされたコードは、分岐命令のときだけプレディケートレジスタを使用する^[要出典]。条件分岐は以下のようになる。

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

setp.cc.type命令は、適切な型の2つのレジスタを比較した結果をプレディケートレジスタに格納する。set命令も存在する。set.le.u32.u64 %r101, %rd12, %rd28というコードは、もしも64bitレジスタ %rd12 が64bitレジスタ %rd28 よりも小さい、あるいは等しいならば、32bitレジスタ %r101 へ 0xffffffff を格納する。さもなければ、%r101は、0x00000000 を格納される。

疑似レジスタを意味する少数の定義済み識別子がある。他にも %tid, %ntid, %ctaid そして %nctaid は、それぞれスレッドのインデックス、ブロックの次元、ブロックのインデックス、そして、グリッドの次元を格納している^[3]。

状態空間

ロード (ld) とストア (st) 命令は、複数の明確な状態空間（メモリバンク）の一つを参照する。例えば、ld.paramの場合、状態空間.paramを参照している。

8つの状態空間が存在する^[3]。

.reg : レジスタ
.sreg : 特別な読み取り専用のプラットフォーム固有のレジスタ
.const : 共有された読み取り専用のメモリ
.global : 全てのスレッドから共有されるグローバルメモリ
.local : それぞれのスレッドだけに属するローカルメモリ
.param : カーネル^[4]に渡すパラメーター
.shared : 1つのブロックの中の複数のスレッドの間で共有されるメモリ
.tex : グローバルテクスチャメモリ（廃止予定）

共有メモリは下記のようにしてPTXファイルの中で宣言される。

.shared .align 8 .b8 pbatch_cache[15744]; // 8バイト境界にアラインされた 15,744 bytes を定義する。

PTXでカーネル関数を書くことは、CUDA Driver APIを経由してPTXモジュールを明示的に登録することを要求する。このことは、CUDA Runtime APIとNVIDIAのCUDAコンパイラであるnvccを使うよりもやっかいなことになることが多い。GPU Ocelot プロジェクトは、CUDA Runtime APIのカーネル呼び出し機能と一緒にPTXモジュールを登録するAPIを提供したが、GPU Ocelot は、もはや積極的にメンテナンスされていない^[5]。

出典

^ "User Guide for NVPTX Back-end — LLVM 7 documentation". llvm.org.
^ コンピュータアーキテクチャの話(350) NVIDIAの世代別GPUに見るハードウェアの違い | マイナビニュース
^ ^a ^b "PTX ISA Version 2.3" (PDF).
^ GPU内に存在する複数のCUDAコア上で並列実行される処理を記述した関数をカーネル (kernel) と呼ぶ。ストリーム・プロセッシングで用いられる用語であり、OpenCLなどでも使われている。
^ "Google Code Archive - Long-term storage for Google Code Project Hosting". code.google.com.

外部リンク

PTX ISA Version 1.4 NVIDIA, 2009-03-31
PTX ISA Version 2.3 NVIDIA, 2011-11-03
PTX ISA Version 3.2 NVIDIA, 2013-07-19
PTX ISA Version 4.0 NVIDIA, 2014-04-12
PTX ISA Version 4.3 NVIDIA, 2015-08-15
PTX ISA Version 5.0 NVIDIA, 2017-06-xx
PTX ISA Version 6.0 NVIDIA, 2017-09-xx
PTX ISA Version 6.3 NVIDIA, 2018-10-xx
PTX ISA Version 6.4 NVIDIA, 2019-02-xx
PTX ISA page on NVIDIA Developer Zone
GPU Ocelot, April 2011

[1] "User Guide for NVPTX Back-end — LLVM 7 documentation". llvm.org.

[2] コンピュータアーキテクチャの話(350) NVIDIAの世代別GPUに見るハードウェアの違い | マイナビニュース

[ptx-isa-3] "PTX ISA Version 2.3" (PDF).

[4] GPU内に存在する複数のCUDAコア上で並列実行される処理を記述した関数をカーネル (kernel) と呼ぶ。ストリーム・プロセッシングで用いられる用語であり、OpenCLなどでも使われている。

[5] "Google Code Archive - Long-term storage for Google Code Project Hosting". code.google.com.

[1]

[2]

[3]

[4]

[5]

Parallel_Thread_Executionとは？わかりやすく解説