CUDA Dynamic Parallelism API

子カーネル関数立ち上げ

CUDA Dynamic Parallelism APIを用いることでCUDAではカーネル関数内から更にカーネル関数を呼ぶことができます．
スレッドの構造はこの図のようになります．

ただし無制約で関数を呼べるわけではありません．

引数アドレス制約

例えば

__global__ void parent_kernel(){
    int a = 0;
    child_kernel<<<1, 128>>>(&a);
    // ...

のようにローカル変数のアドレスを子カーネル関数に渡すことはできません． この場合は

__device__ int a;
__global__ void parent_kernel(){
    child_kernel<<<1, 128>>>(&a);
    // ...

とするとコンパイルできます．

PTXで見るCUDA Dynamic Parallelism API

CUDAの中間アセンブリであるPTXで子カーネルの立ち上げを見てみます．

子カーネル立ち上げのcuコードからptxを出力すると.weakディレクティブ付きでcudaLaunchDeviceV2関数やcudaGetDevice関数、cudaMalloc関数などのダミー関数がPTX内に定義されています．
.weakですのでリンク時により強い宣言がある場合上書きされます．(そりゃdevice側のcudaMallocがあったとして，そのPTXをNVIDIAが簡単に我々下々の開発者に公開するわけがないですよね)

肝心の子カーネル関数の立ち上げ部分の流れは次のようになります．

1. cudaGetParameterBufferV2関数で子カーネル関数のアドレス，ブロック数，スレッド数などを指定します．
2. cudaLaunchDeviceV2関数で子カーネル関数の引数を格納した領域のアドレスなどを渡し，子カーネル関数を立ち上げます．

実験

デバイス側からのカーネル関数立ち上げが高速だった場合，<<<1, 1>>>で立ち上げ用のカーネル関数を立ち上げ，内部で目的のカーネル関数を立ち上げた方が速く立ち上げるかもと思い試してみましたが遅かったです．

実験コード : gitlab.momo86.net : mutsuki/cuda-child-kernel

あとがき

子カーネル関数の呼び出し部分のPTXを読んでいたらretvalなどという若干意味の伝わない変数名が用いられていて，リーダブルコードでも読めといいたくなりました．

参考文献

• NVIDIA Developer blog - CUDA Dynamic Parallelism API and Principles