消えたCUDA関連の旧ブログ記事を復元するひとり Advent Calendar 2024の記事です。

Pascal世代から使えるようになった半精度演算はPTXでの型としてはf16とf16x2にあたる. f16x2はf16を2つまとめたもので,これを使うとSIMD演算をすることができる. 今回は備忘録も兼ねてf16x2のことを少しだけ書いておく.

f16x2

レジスタ

.reg .b32 %hoge;
のように32bit型で指定する.

計算

mul.f16x2 %hoge0, %hoge1, %hoge2;
のようにf16x2型で計算する.

f16,f16x2の相互変換

変換というほどでもないが 
.reg .b16 %splited<2>;

//f16x2を2つのf16へ
mov.b32 { %splited0, %splited1 }, %hoge;
//2つのf16をf16x2へ
mov.b32 %hoge, { %splited0, %splited1 };

組み込み関数

PTXは書きたくない!という方はCUDAの組み込み関数を使えばf16x2の演算を使えます.(こっちが普通の使い方かも?)

参考ページ