【CUDA】Dynamic Parallelism

Dynamic parallelismとは、カーネルの中からカーネルを呼び出すことができるsm_35以降の機能です。

ちょっと具体的にどういった場面に使えるかよくわからないのですが、
Adaptive Parallel Computation with CUDA Dynamic Parallelism | Parallel Forall
によると、

• adaptive gridsのような階層的なデータ構造を用いるアルゴリズム
• クイックソートのような再帰を用いるアルゴリズム

などが挙げられています。

#include <stdio.h>

__global__ void ChildKernel(void* data){

	printf("child : %d, %d\n", blockIdx.x, threadIdx.x);

}

__global__ void ParentKernel(void* data){

	printf("parent: %d, %d\n", blockIdx.x, threadIdx.x);

	ChildKernel<<<1, 2>>>(data);
	cudaDeviceSynchronize();

}

int main(int argc, char** argv){
	int nbytes = 1024;
	int *data_dev = 0;
	cudaMalloc((void**)&data_dev, nbytes);
	cudaMemset(data_dev, 255, nbytes);


	ParentKernel<<<1, 2>>>(data_dev);
	return 0;
}

コンパイル時には，-arch=sm_35と-rdc=trueの指定が必要