Dynamic Parallelism 机制实现两个 cuda kernel 动态并行的示例

1，简介 Dynamic Parallelism

Dynamic Parallelism 是 CUDA 的一项特性，允许一个 CUDA kernel 在 GPU 上直接启动其他 CUDA kernel，而无需返回主机。这种特性使得在 GPU 上可以实现更复杂的并行盘算模式，特别实用于递归算法、分条理盘算和自顺应网格细化等场景。

2，示例一：kernel 启动单个 kernel

2. #include <iostream>
3. #include <cuda_runtime.h>
5. // 子 kernel，计算每个元素的平方
6. __global__ void childKernel(int* data, int n) {
7.     int idx = blockIdx.x * blockDim.x + threadIdx.x;
8.     if (idx < n) {
9.         data[idx] +=100;
10.     }
11. }
13. // 父 kernel，启动子 kernel
14. __global__ void parentKernel(int* data, int n) {
15.     // 假设每个 block 处理 256 个元素
16.     int threadsPerBlock = 1024;
17.     int blocks = (n + threadsPerBlock - 1) / threadsPerBlock;
19.     // 启动子 kernel
20.     childKernel<<<blocks, threadsPerBlock>>>(data, n);
22.     // 等待子 kernel 完成
23. //    cudaDeviceSynchronize();
24. }
26. // 主机代码
27. int main() {
28.     const int n = 1024*1024*512;
29.     int* h_data = new int[n];
30.     int* d_data;
32.     // 初始化数据
33.     for (int i = 0; i < n; ++i) {
34.         h_data[i] = i;
35.     }
37.     // 分配设备内存
38.     cudaMalloc(&d_data, n * sizeof(int));
39.     cudaMemcpy(d_data, h_data, n * sizeof(int), cudaMemcpyHostToDevice);
41.     // 启动父 kernel
42.     parentKernel<<<1, 1>>>(d_data, n);
44.     // 复制结果回主机
45.     cudaMemcpy(h_data, d_data, n * sizeof(int), cudaMemcpyDeviceToHost);
47.     // 输出结果
48.     for (int i = 0; i < 10; ++i) {
49.         std::cout << h_data[i] << " ";
50.     }
51.     std::cout << std::endl;
53.     // 清理
54.     delete[] h_data;
55.     cudaFree(d_data);
57.     return 0;
58. }

复制代码

编译：
 nvcc hello.cu -rdc=true -arch=sm_75 -lcudart -Xcompiler -lcudadevrt
大概Makefile：

3. DEV_FLAGS :=  -rdc=true -arch=sm_75 -Xcompiler
5. EXE := hello
7. aLL: $(EXE)
9. %: %.cu
10.         nvcc $< -o $@ $(DEV_FLAGS) -lcudart -lcudadevrt
12. .PHONY: clean
13. clean:
14.         -rm -rf $(EXE)

复制代码

运行：

 
3，示例二：kernel 启动两个 kernel

2. #include <iostream>
3. #include <cuda_runtime.h>
5. // data[2k] += 100;
6. __global__ void childKernel_A(int* data, int n) {
7.     int idx = blockIdx.x * blockDim.x + threadIdx.x;
8.     if (idx < n/2) {
9.         data[2*idx] += 100;
10.     }
11. }
13. //data[2k+]  += 200;
14. __global__ void childKernel_B(int* data, int n) {
15.     int idx = blockIdx.x * blockDim.x + threadIdx.x;
16.     if (idx < n/2) {
17.         data[2*idx+1] += 200;
18.     }
19. }
20. // launch kernels
21. __global__ void parentKernel(int* data, int n) {
23.         if(threadIdx.x==0){
24.                 int threadsPerBlock = 1024;
25.                 int blocks = (n/2 + threadsPerBlock - 1) / threadsPerBlock;
27.                 childKernel_A<<<blocks, threadsPerBlock>>>(data, n);
28.                  //cudaDeviceSynchronize();
29.         }
30.         else if(threadIdx.x==1){
31.                
32.                 int threadsPerBlock = 1024;
33.                 int blocks = (n/2 + threadsPerBlock - 1) / threadsPerBlock;
35.                 childKernel_B<<<blocks, threadsPerBlock>>>(data, n);
36.         }
37. }
39. int main() {
40.     const int n = 1024*1024*512;
41.     int* h_data = new int[n];
42.     int* d_data;
44.     for (int i = 0; i < n; ++i) {
45.         h_data[i] = i;
46.     }
48.     cudaMalloc(&d_data, n * sizeof(int));
49.     cudaMemcpy(d_data, h_data, n * sizeof(int), cudaMemcpyHostToDevice);
51.     parentKernel<<<1, 2>>>(d_data, n);
53.     cudaMemcpy(h_data, d_data, n * sizeof(int), cudaMemcpyDeviceToHost);
55.    
56.     for (int i = 0; i < 10; ++i) {
57.         std::cout << h_data[i] << " ";
58.     }
59.     std::cout << std::endl;
61.     delete[] h_data;
62.     cudaFree(d_data);
64.     return 0;
65. }

复制代码

 nvcc hello.cu -rdc=true -arch=sm_75 -lcudart -Xcompiler -lcudadevrt
运行：

4，总结：


4.1 上风

镌汰主机干预: 通过在装备上直接启动子 kernel，镌汰了主机与装备之间的通讯开销。
简化代码结构: 允许在 GPU 上实现更复杂的控制流和盘算模式。
进步性能: 在某些场景下，通过镌汰主机干预和优化并行盘算，可以进步团体性能。
4.2 利用场景

须要与其他 GPU 盘算框架集成，须要风雅控制 GPU 资源和举动。
递归算法: 如快速排序、树形数据结构遍历。
自顺应网格细化: 在盘算流体力学和其他科学盘算中。
分条理盘算: 如多级网格方法。