CUDAでのKernel関数のポインタテーブル

CUDAでは最適化のテクニックとしてKernel関数をC++のtemplateを利用して量産して、適切な関数を呼び出すというのが結構使える。
たとえばKernel関数内部の最内のループ回数が1〜10の間で可変のときに、
ループする回数ごとに10種類の関数を用意して利用する。
ループ回数が決まっていればループのアンロールが可能になるので、速度が向上する可能性がある。

（templateを使わずに、コピペで大量の関数を用意することもできるけど保守性を考えたら絶対にお勧めしない。）

関数が大量になってきたときにswitch分とかで関数を選んでいると、コードがどんどん長くなるので、__global__の関数はポインタがあることを利用して関数ポインタのテーブルを利用すると良い。

次のような感じになる。

// templateを使った関数の簡単な例
template  void kernel_func(int *dest){
    int i, j;
    for (i = 0; i < a; i++){
        for (j = 0; j < b; j++){
             dest[threadIdx.x] += i * j;
        }
    }
}

// aの値とbの値がそれぞれ0〜1の値をとる場合
// 関数ポインタのテーブル
static void (* const kernel_tbl[3][3])(int *) = {
    {
        kernel_func<0, 0>,
        kernel_func<0, 1>,
    },
    {
        kernel_func<1, 0>,
        kernel_func<1, 1>,
    },
}

呼び出しはこんな感じ

kernel_tbl[index_a][index_b]<<>>(d_Out);