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

例えばカーネル関数内で2冪の整数のlog2を取りたいときどうしますか?
浮動小数点数に変換してlog2をとったり 

__device__ unsigned log2lu(const unsigned long a) {
    unsigned i = 0;
    for (; (1lu << (i + 1)) <= a; i++);
    return i;
}
みたいなコードを書いたりするかもですよね．
あまり知られていないかもしれませんが，CUDAには整数 & bit演算用の組み込み関数みたいなものがあり，これを使うと 
__device__ unsigned long log2lu(const unsigned long a) {
    return 64 - __ffsll(__brevll(a));
}
とforやifなしに書くことができます（a=0は気にしない）．

こういう関数はどういう種類があるのかを知っておかないといざというときに使えないのでまとめておきます．

整数 & bit演算 関数

bit操作/カウント

32bit変数用関数 64bit変数用関数 動作
__brev
 __brevll bit列を反転
__clz __clzll 最上位ビットから探索を始め，連続する0の個数
__ffs __ffsll 最下位bitの位置
- __popcll 64bitのうち1の立っている個数

__byte_perm(x, y, s)

言葉で説明するのが大変な関数です．
簡単に言うと2つのu32変数x,yから一部を切り出して新しい変数を作る関数です．
切り出し方をu32変数引数sで指定します．
図を見てなんとなくわかってもらえれば．

funnelshift関数family

4つあります．

  • __funnelshift_l (lo, hi, shift) : 32uを2つ連結{hi : lo}しu64とした状態で(shift & 0xff) bit左シフトし，上位32bitを返す
  • __funnelshift_lc (lo, hi, shift) : 32uを2つ連結{hi : lo}しu64とした状態でmin(shift, 32) bit左シフトし，上位32bitを返す
  • __funnelshift_r (lo, hi, shift) : 32uを2つ連結{hi : lo}しu64とした状態で(shift & 0xff) bit右シフトし，下位32bitを返す
  • __funnelshift_rc (lo, hi, shift) : 32uを2つ連結{hi : lo}しu64とした状態でmin(shift, 32) bit右シフトし，下位32bitを返す

    • 癖のある算術関数っぽいもの

    Signed変数引数関数 Unsigend変数引数関数 動作
    __hadd(a_s32, b_s32)
    		__uhadd(a_u32, b_u32) (a + b) / 2 を加算部分でのオーバーフローを避けつつ計算（切り捨て）
    __rhadd(a_s32, b_s32) __urhadd(a_u32, b_u32) (a + b) / 2 を加算部分でのオーバーフローを避けつつ計算（切り上げ）
    __mul24(a_s32, b_s32) __umul24(a_u32, b_u32) 積a[23:0] * b[23:0]計算
    __mul64hi(a_s64, b_s64) __umul64hi(a_u64, b_u64) 積a * bの結果128bitのうち上位64bitを返す
    __mulhi(a_s32, b_s32) __umulhi(a_u32, b_u32) 積a * bの結果64bitのうち上位32bitを返す
    __sad(a_s32, b_s32, z_u32) __usad(a_u32, b_u32, z_u32) 差の絶対値を加算 |a - b| + z

    おわり

    これらの関数はそのままPTXで同じような名前の命令になるものもあれば，他いくつかの命令によって構成されるものもあるようです．

    参考

    CUDA Math Function, 1.7. Integer Intrinsics