CUDA旧动态并行弃用后，CDP2动态并行同步的新最佳实践

直到几年前，以下演示性CUDA代码仍是完全可行的：

__global__ void parent_kernel(int *output){

  int tid = blockIdx.x * blockDim.x + threadIdx.x;
  int Alice, Bob;
  float *Charly, *Dave;

  Alice = tid;
  for(int j=0; j<100; ++j){
    child_kernel1<<<1,512>>>(Charly);
    cudaThreadSynchronize();
    Alice=f(Alice,Charly);

    child_kernel2<<<1,512>>>(Dave);
    cudaThreadSynchronize();
    Bob=g(Alice,Dave);
  }
  output[tid] = Bob;
}


//Later...
parent_kernel<<<32,32>>>(Ellie);

这段代码的核心限制在于：for循环无法并行化，因为迭代j+1依赖于j的计算结果，且调用的f、g函数可能存在强非线性特性。原方案通过设备端调用子核（动态并行CDP）处理大型数组，父核线程需等待子核所有线程执行完成后，才能继续计算f和g——该方法曾出自NVIDIA官方文档，但目前已被CUDA编程指南标记为弃用，且无零成本替代方案，必须重构代码。

以下是基于最新CDP2框架的最佳实践，同时兼顾实际场景中可能存在的多函数、嵌套if语句及动态子核线程数需求：

核心重构思路 #

将原父核的启动配置改为：

parent_kernel<<<1024,512>>>(Ellie);

移除父核内的子核调用，直接利用块内的512个线程完成原属于子核的计算任务。

关于Alice=f(Alice,Charly)的执行逻辑 #

针对这个问题，需结合数据依赖和线程分工明确：

• 线程分工与数据作用域：
  如果Alice是每个父线程独有的私有变量（如原代码中每个线程的tid初始化值），则每个线程需独立执行Alice=f(Alice,Charly)；若Charly是子核计算产生的块级共享数据，必须通过共享内存或全局内存完成传递：重构后由父核块内的512个线程协作完成Charly的计算，之后通过__syncthreads()同步块内所有线程，确保所有线程都能读取到完整的Charly数据后，再各自执行f函数更新私有变量Alice。
• 同步与数据一致性：
  执行f函数前，必须通过__syncthreads()确保块内所有线程完成Charly的计算，避免数据未就绪；若存在嵌套if语句，需注意线程发散问题，可通过条件同步或掩码同步（如__syncwarp()）保证数据一致性。

额外优化建议 #

• 针对动态子核线程数需求：可通过块内线程的动态分工（如根据条件分配线程任务）替代原动态子核的启动逻辑，利用threadIdx.x进行任务划分。
• 多函数嵌套场景：将原各子核的计算逻辑封装为__device__函数，在父核块内线程中直接调用，减少核启动开销的同时，保持代码的模块化。

内容的提问来源于stack exchange，提问作者Valentin Aslanyan