CUDA PTX

PTXとは

CUDAの中間アセンブリ言語と言ったところでしょうか. さらに機械語に近いSASS(NVIDIAの人はサスと読んでいた)というアセンブリもある.

特徴

• 人間が読める
• 手軽に自分のソースコードに組み込める
• 公式のリファレンスが一応ある

PTXを読む

バイナリ内のPTXを出力

# cuobjdump -ptx filename

でfilenameに含まれているPTXコードを読むことができます.

勘違いしがちですが、このコマンドは機械語からPTXを生成しているのではなく、バイナリ内に含まれているPTXを出力しています. そのため

# cuobjdump -sass filename

で逆アセンブルをして出力されるSASSと内容が異なる場合があります.

.cuコードからPTXを出力

# nvcc -ptx filename.cu

読み方

読んでみれば分かりますが命令がとても分かりやすいので次のことくらいを知っていればいいと思います.

• 基本は

  演算名.型名 レジスタ レジスタ レジスタ
  

  の構成になっている.

  演算名.オプション.型名.型名 レジスタ レジスタ レジスタ
  

  のように命令名の後に丸め方や入力型,出力型をたくさん書くものもある.
• レジスタは自分で型を選べ、名前をつけられる

PTXを書く

インラインアセンブラを使って自分のコードに簡単に埋め込むことができます.

fp16計算で16bit浮動小数2つに対して同時に2乗計算を行うプログラム

C++11以降の生文字列リテラルを使うとスッキリ書けます.

__global__ void square_f16x2(half2* a,int n){
	asm(R"(
{
	.reg .u32 %mr<4>;
	.reg .u64 %memory_at;
	.reg .b32 %memory_val;
	.reg .b32 %memory_res;

	mov.s32 %mr0, %ctaid.x;
	mov.s32 %mr1, %ntid.x;
	mov.s32 %mr2, %tid.x;
	mad.lo.s32 %mr3, %mr0, %mr1, %mr2;

	mad.wide.s32 %memory_at, %mr3, 4, %0;
	ld.global.b32 %memory_val, [%memory_at];
	mul.f16x2 %memory_res , %memory_val, %memory_val;
	st.global.b32 [%memory_at], %memory_res;
}
)"
	::"l"(a),"r"(n)
	);
}

終わりに

自分でPTXでコードを書いたほうがnvccの最適化より速くなったこともあるので，極めれば...

参考ページ

• NVIDIA PTX ISA
• NVIDIA Inline Assembly