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

何がしたいか

CUDAにはstreamと呼ばれるスケジューリング機能があり、これを用いることでメモリの非同期コピーなどができます。
CUDA 10からはCUDA Graphと呼ばれるstreamをグラフとして記述し実行する機能もあり、複数のカーネル関数を1つ1つ立ち上げるのと比較してGraphとして立ち上げたほうが高速に立ち上げられるという結果もあります[1]。
で、今回はそんなstreamにCPU側の処理を流したい場合にどうすればいいかという話です。

サンプルコード

cuLaunchHostFunc関数[2]を使うと流すことができます。 

#include <cstdio>
#include <cuda.h>

struct data_struct_t {
  int a, b, c;
};

void host_func(void* const data_void_ptr) {
  auto* data_ptr = reinterpret_cast<data_struct_t*>(data_void_ptr);

  data_ptr->a = 10;
  data_ptr->b = 20;
  data_ptr->c = 30;
}

int main() {
  cudaStream_t cuda_stream;
  cudaStreamCreate(&cuda_stream);

  data_struct_t data;

  cuLaunchHostFunc(cuda_stream, &host_func, reinterpret_cast<void*>(&data));

  cudaStreamSynchronize(cuda_stream);

  std::printf("a=%d, b=%d, c=%d\n", data.a, data.b, data.c);
}

コンパイル時には 

nvcc main.cu -std=c++11 -arch=sm_70 -lcuda
の様にlibcudaをリンクさせる必要があります。
cuLaunchHostFunc関数に渡すhost側の関数の引数はvoid*とし、cuLaunchHostFunc関数の第３引数に渡します。
気持ちとしてはWindows APIのCreateThreadの気持ちです（伝われ）。

cuLaunchHostFunc関数はExecution Controlの一部ですが、CUDA 10.1でAPIが刷新されており、これによって追加された関数です。
そのためCUDA 10.0以前では使えないのでお気をつけください。 また、制約として、cudaMallocなどのCUDA APIは、引数で渡す関数内で実行できません。

参考

  1. NVIDIA A100 Tensor Core GPU Architecture, NVIDIA, (p60 : Task Graph Acceleration on NVIDIA Ampere Architecture GPUs)
  2. cuLaunchHostFunc - Execution Control, CUDA TOOLKIT DOCUMENTATION