如何确保CUDA子核完成处理后父核再继续？动态并行求和问题

正在学习CUDA，制作基础图像盒式模糊Demo时遇到问题：用子核sumValues计算邻域像素和，结果始终为0。排查发现子核未完成处理父核就继续执行，且启动的核数量超出设备限制。之前知道可以在设备代码中用cudaDeviceSynchronize()，但该方法已废弃，想知道如何实现子核完成后父核再执行，或者用动态并行计算邻域和的思路本身是否有误？

附上代码：

__global__ void sumValues(uchar3* img, Pair2i* size, int origin, Pair3u* tmpSums, uchar kernel_size)
{   
    int Idx = origin + (threadIdx.y - (int)ceil((kernel_size - 1) / 2.0)) * size->y + (threadIdx.x - (int)ceil((kernel_size - 1) / 2.0));

    if (Idx < 0 || Idx >= size->x * size->y) return;

    atomicAdd((&tmpSums[origin].x), unsigned int(img[Idx].x));
    atomicAdd((&tmpSums[origin].y), unsigned int(img[Idx].y));
    atomicAdd((&tmpSums[origin].z), unsigned int(img[Idx].z));
}

__global__ void blur(uchar3* img, Pair2i* size, uchar3* res, Pair3u* tmpSums, uchar kernel_size)
{
    unsigned int Idx = access2d(threadIdx, blockIdx, blockDim, gridDim);

    if (Idx >= size->x * size->y) return;

    tmpSums[Idx] = { 0, 0, 0 };

    dim3 dimBlock(kernel_size, kernel_size);

    sumValues<<<1, dimBlock>>>(img, size, Idx, tmpSums, kernel_size);

    res[Idx].x = tmpSums[Idx].x / (kernel_size * kernel_size);
    res[Idx].y = tmpSums[Idx].y / (kernel_size * kernel_size);
    res[Idx].z = tmpSums[Idx].z / (kernel_size * kernel_size);
}

补充说明：使用OpenCV读取图像，将数据指针传递给计算模糊图像的核函数。父核blur调用子核sumValues计算求和结果并存储到tmpSums中，随后父核计算均值并赋值给结果图像。

1. 子核与父核的同步处理 #

设备端动态并行中，父核启动子核后不会自动等待其完成，这是结果为0的直接原因。注意：废弃的是主机端部分场景下的cudaDeviceSynchronize()用法，设备端调用该函数仍然有效，可在父核中sumValues启动后添加：

cudaDeviceSynchronize();

这会强制父核线程等待所有子核执行完毕，再继续计算均值。

2. 核数量超限问题的本质 #

当前每个图像像素对应一个父核线程，每个父核又启动一个子核，这会生成与图像像素数等量的子核（比如1080P图像就有200多万个子核），远远超过CUDA设备的并发上限（通常仅支持数百个活跃核）。这种做法会导致大量子核被阻塞甚至无法启动，最终求和结果为0，同时带来巨大的调度开销。

3. 动态并行思路的合理性判断 #

用动态并行实现盒式模糊的邻域求和完全不合理：动态并行的设计目标是处理递归或嵌套并行的复杂任务，而盒式模糊的邻域计算属于局部、规则的并行任务，用动态并行属于过度设计，性能反而远低于常规优化方案。

推荐替代方案：共享内存优化 #

在单个核函数内利用共享内存缓存邻域像素，避免重复读取全局内存，直接计算邻域和，是盒式模糊的标准优化方式：

__global__ void blur_shared(uchar3* img, Pair2i* size, uchar3* res, uchar kernel_size)
{
    const int tx = threadIdx.x;
    const int ty = threadIdx.y;
    const int bx = blockIdx.x;
    const int by = blockIdx.y;
    const int block_w = blockDim.x;
    const int block_h = blockDim.y;

    const int x = bx * block_w + tx;
    const int y = by * block_h + ty;
    const int global_idx = y * size->x + x;

    if (x >= size->x || y >= size->y) return;

    const int half_k = (kernel_size - 1) / 2;
    const int shared_w = block_w + 2 * half_k;
    extern __shared__ uchar3 shared_img[];

    // 加载共享内存（含边界扩展）
    const int s_x = tx + half_k;
    const int s_y = ty + half_k;
    shared_img[s_y * shared_w + s_x] = img[global_idx];

    // 左边界填充
    if (tx < half_k) {
        const int left_x = max(bx * block_w + tx - half_k, 0);
        shared_img[s_y * shared_w + tx] = img[y * size->x + left_x];
    }
    // 右边界填充
    if (tx >= block_w - half_k) {
        const int right_x = min(bx * block_w + tx + half_k, size->x - 1);
        shared_img[s_y * shared_w + s_x + half_k] = img[y * size->x + right_x];
    }
    // 上边界填充
    if (ty < half_k) {
        const int top_y = max(by * block_h + ty - half_k, 0);
        shared_img[ty * shared_w + s_x] = img[top_y * size->x + x];
    }
    // 下边界填充
    if (ty >= block_h - half_k) {
        const int bottom_y = min(by * block_h + ty + half_k, size->y - 1);
        shared_img[(s_y + half_k) * shared_w + s_x] = img[bottom_y * size->x + x];
    }

    __syncthreads();

    // 计算邻域和
    unsigned int sum_x = 0, sum_y = 0, sum_z = 0;
    for (int dy = -half_k; dy <= half_k; dy++) {
        for (int dx = -half_k; dx <= half_k; dx++) {
            const uchar3 pixel = shared_img[(s_y + dy) * shared_w + (s_x + dx)];
            sum_x += pixel.x;
            sum_y += pixel.y;
            sum_z += pixel.z;
        }
    }

    const int kernel_area = kernel_size * kernel_size;
    res[global_idx].x = sum_x / kernel_area;
    res[global_idx].y = sum_y / kernel_area;
    res[global_idx].z = sum_z / kernel_area;
}

大核尺寸最优方案：积分图（前缀和） #

当核尺寸较大时，积分图方案效率更高：先计算图像的积分图，每个像素的邻域和可通过积分图的四个顶点值快速推导，时间复杂度与核尺寸无关，仅为O(W×H)。

总结 #

• 同步问题可通过在父核中调用设备端cudaDeviceSynchronize()解决，但因核数量超限，该方案不具备实用价值。
• 动态并行不适合盒式模糊这类局部规则计算任务，建议改用共享内存优化或积分图方案，既能解决结果异常问题，又能大幅提升计算效率。

内容的提问来源于stack exchange，提问作者Amr Emad