読者です 読者をやめる 読者になる 読者になる

【CUDA】Dynamic Parallelism

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

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

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

#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の指定が必要