「Parallel Thread Execution」の版間の差分
m カテゴリ追加 |
|||
3行目: | 3行目: | ||
== レジスタ == |
== レジスタ == |
||
PTXは、任意の大きさの[[レジスタ (コンピュータ)#レジスタセット|レジスタセット]]を使用する。コンパイラからの出力は、ほぼ純粋な単一代入形式である。通常、連続したレジスタを参照する連続した命令列を伴っている。プログラムは以下の形式の宣言で始まる。 |
PTXは、任意の大きさの[[レジスタ (コンピュータ)#レジスタセット|レジスタセット]]を使用する。コンパイラからの出力は、ほぼ純粋な単一代入形式である。通常、連続したレジスタを参照する連続した命令列を伴っている。プログラムは以下の形式の宣言で始まる。 |
||
< |
<syntaxhighlight lang="asm"> |
||
.reg .u32 %r<335>; // 符号なし32bit整数型の335個のレジスタを宣言する。 |
.reg .u32 %r<335>; // 符号なし32bit整数型の335個のレジスタを宣言する。 |
||
</syntaxhighlight> |
|||
</source> |
|||
PTXは、最大で3つの引数を持つアセンブリ言語である。そして、ほとんど全ての命令は、命令が操作するデータのデータ型を(符号とデータ幅によって)明示するために命令の右にデータ型を並べて書く。レジスタ名は、先頭に%文字を付ける。そして、定数はそのまま書く。以下に例を示す。 |
PTXは、最大で3つの引数を持つアセンブリ言語である。そして、ほとんど全ての命令は、命令が操作するデータのデータ型を(符号とデータ幅によって)明示するために命令の右にデータ型を並べて書く。レジスタ名は、先頭に%文字を付ける。そして、定数はそのまま書く。以下に例を示す。 |
||
< |
<syntaxhighlight lang="asm"> |
||
shr.u64 %rd14, %rd12, 32; // %rd12の符号なし64bit整数を32bit右シフトする。結果は%rd14に入る。 |
shr.u64 %rd14, %rd12, 32; // %rd12の符号なし64bit整数を32bit右シフトする。結果は%rd14に入る。 |
||
cvt.u64.u32 %rd142, %r112; // 符号なし32bit整数を64bitへ変換する。 |
cvt.u64.u32 %rd142, %r112; // 符号なし32bit整数を64bitへ変換する。 |
||
</syntaxhighlight> |
|||
</source> |
|||
プレディケートレジスタもある(命令を実行するかどうかを決める条件を格納するレジスタ)。しかし、{{要出典範囲|date=2019-06|シェーダーモデル1.0でコンパイルされたコードは、分岐命令のときだけプレディケートレジスタを使用する}}。条件分岐は以下のようになる。 |
プレディケートレジスタもある(命令を実行するかどうかを決める条件を格納するレジスタ)。しかし、{{要出典範囲|date=2019-06|シェーダーモデル1.0でコンパイルされたコードは、分岐命令のときだけプレディケートレジスタを使用する}}。条件分岐は以下のようになる。 |
||
< |
<syntaxhighlight lang="asm"> |
||
@%p14 bra $label; // $labelへ分岐 |
@%p14 bra $label; // $labelへ分岐 |
||
</syntaxhighlight> |
|||
</source> |
|||
<code>setp.cc.type</code>命令は、適切な型の2つのレジスタを比較した結果をプレディケートレジスタに格納する。<code>set</code>命令も存在する。< |
<code>setp.cc.type</code>命令は、適切な型の2つのレジスタを比較した結果をプレディケートレジスタに格納する。<code>set</code>命令も存在する。<syntaxhighlight lang="asm" inline>set.le.u32.u64 %r101, %rd12, %rd28</syntaxhighlight>というコードは、もしも64bitレジスタ <code>%rd12</code> が64bitレジスタ <code>%rd28</code> よりも小さい、あるいは等しいならば、32bitレジスタ <code>%r101</code> へ <code>0xffffffff</code> を格納する。さもなければ、<code>%r101</code>は、<code>0x00000000</code> を格納される。 |
||
疑似レジスタを意味する少数の定義済み識別子がある。他にも <code>%tid, %ntid, %ctaid</code> そして <code>%nctaid</code> は、それぞれスレッドのインデックス、ブロックの次元、ブロックのインデックス、そして、グリッドの次元を格納している<ref name="ptx-isa">{{cite web2|url=http://developer.download.nvidia.com/compute/cuda/3_1/toolkit/docs/ptx_isa_2.3.pdf|title=PTX ISA Version 2.3|publisher=}}</ref>。 |
疑似レジスタを意味する少数の定義済み識別子がある。他にも <code>%tid, %ntid, %ctaid</code> そして <code>%nctaid</code> は、それぞれスレッドのインデックス、ブロックの次元、ブロックのインデックス、そして、グリッドの次元を格納している<ref name="ptx-isa">{{cite web2|url=http://developer.download.nvidia.com/compute/cuda/3_1/toolkit/docs/ptx_isa_2.3.pdf|title=PTX ISA Version 2.3|publisher=}}</ref>。 |
||
37行目: | 37行目: | ||
共有メモリは下記のようにしてPTXファイルの中で宣言される。 |
共有メモリは下記のようにしてPTXファイルの中で宣言される。 |
||
< |
<syntaxhighlight lang="nasm"> |
||
.shared .align 8 .b8 pbatch_cache[15744]; // 8バイト境界にアラインされた 15,744 bytes を定義する。 |
.shared .align 8 .b8 pbatch_cache[15744]; // 8バイト境界にアラインされた 15,744 bytes を定義する。 |
||
</syntaxhighlight> |
|||
</source> |
|||
<!-- mov.u64 %rd9, pbatch_cache; |
<!-- mov.u64 %rd9, pbatch_cache; |
||
Shared memory is generally addressed via a kernel-global pointer set up at the start of the kernel by |
Shared memory is generally addressed via a kernel-global pointer set up at the start of the kernel by |
2020年7月5日 (日) 23:11時点における版
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.
{{cite web}}
: Cite webテンプレートでは|access-date=
引数が必須です。 (説明) - ^ コンピュータアーキテクチャの話(350) NVIDIAの世代別GPUに見るハードウェアの違い | マイナビニュース
- ^ a b "PTX ISA Version 2.3" (PDF).
{{cite web}}
: Cite webテンプレートでは|access-date=
引数が必須です。 (説明) - ^ GPU内に存在する複数のCUDAコア上で並列実行される処理を記述した関数をカーネル (kernel) と呼ぶ。ストリーム・プロセッシングで用いられる用語であり、OpenCLなどでも使われている。
- ^ "Google Code Archive - Long-term storage for Google Code Project Hosting". code.google.com.
{{cite web}}
: Cite webテンプレートでは|access-date=
引数が必須です。 (説明)
外部リンク
- 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