CUDA动态并行性能较普通内核低30倍的原因及实现疑问

注：当前未携带电脑与GPU，以下内容凭记忆整理，相关代码已完成编译与计时验证，若存在零散笔误可忽略。

我无法确认观测到的性能问题是机制开销导致、还是自身实现有误，也不明确为何内核内启动子内核的CUDA动态并行模式，比采用线程谓词判断、存在部分线程闲置的单一大内核运行速度更慢，初步推测可能是当前测试任务量过小，未达到GPU工作饱和状态，无法体现动态并行的优势。

本次测试采用的简单计算任务为：将方阵所有元素乘以2，矩阵尺寸最大不超过16×16。假设设备内存中已存储200个待处理矩阵，首先测试基础内核实现方案的性能。

// One matrix given to each block
__global__ void matrixFunc(Matrix** matrices)
{
    Matrix* m = matrices[blockIdx.x];
    int area = m->width * m->height;
    if (threadIdx.x < area)
        // Heavy calculations
}

// Assume 200 matrices, no larger than 16x16
matrixFunc<<<200, 256>>>(ptrs);

该实现为每个矩阵分配1个线程块，每个块配置充足线程，保证块内线程数不小于矩阵元素总数，实测运行耗时为0.17微秒。

我最初认为该实现存在资源浪费：由于测试用矩阵尺寸普遍很小（例如2×2矩阵仅需4个线程，配置256线程存在大量冗余），因此出于学习验证目的，尝试通过动态并行从父内核内部启动子内核处理各矩阵，测试该模式的运行时开销，对应实现代码如下：

__device__ void matrixFunc(float* matrix)
{
    // Heavy calculations (on threadIdx.x for the cell)
}

__global__ void matrixFuncCaller(Matrix** matrices)
{
    Matrix* m = matrices[threadIdx.x];
    int area = m->width * m->height;
    matrixFunc<<<1, area>>>(m.data);
}

matrixFuncCaller<<<1, 200>>>(ptrs);

该实现性能大幅下降，实测耗时达到11.3微秒。

随后尝试为每个子内核启动配置独立CUDA流进行优化，对应实现逻辑如下：

__global__ void matrixFuncCaller(Matrix** matrices)
{
    Matrix* m = matrices[threadIdx.x];
    int area = m->width * m->height;
    // Create `stream`
    matrixFunc<<<1, area, 0, stream>>>(m.data);
    // Destroy `stream`
}

优化后耗时降至3微秒，性能有所提升，但仍远高于基础实现的0.17微秒，整体性能差距最高达30倍。

核心疑问 #

• 该性能差距是否由动态并行的内核启动开销导致？当前示例任务计算量极小，启动开销可能远高于实际计算耗时；在不便公开的实际业务场景中，单任务计算量比矩阵乘2大很多，但整体任务粒度仍偏小，可能仍存在明显的开销影响。
• 上述动态并行实现方式是否存在错误？

1. 性能差距核心来源是CUDA动态并行（CDP）的固有启动开销 #

你观测到的性能差完全符合动态并行的机制特性：

• 设备端发起内核启动的开销远高于主机端。主机端发起小内核启动的开销通常在1微秒以内，而设备端子内核启动的固有开销即使在新版本CUDA中也在1微秒上下，200个子内核即使做了流并行，累计开销也远高于主机端一次性发起200个线程块的成本。
• 无流版本11.3微秒的高耗时，是因为默认流会强制所有子内核串行执行，200次启动开销直接累加；配置独立流后子内核可以并行调度启动，耗时自然下降，但单次设备端启动的固有成本依然存在，因此无法追上基础实现的性能。
• 你认为的"256线程跑小矩阵存在资源浪费"是对GPU调度逻辑的误解：GPU线程调度是以warp（32线程）为最小单位的，同warp内的线程是锁步执行的。就算是2×2矩阵只用到4个线程，同warp内剩余28个线程通过谓词判断跳过计算的闲置成本几乎可以忽略，远低于额外启动内核的开销。
• 你的测试场景完全不适合用动态并行。动态并行的设计目标是处理主机端无法预知调度逻辑的不规则任务、递归类任务（比如非结构化网格计算、树结构遍历），要求单个子任务的计算量足够大，才能摊薄启动开销。对于已知任务规模、单任务计算量极小的批量处理场景，主机端一次性发起内核、块内用谓词过滤无效线程的方案永远是最优解。

2. 你的动态并行实现无逻辑错误，但存在可优化点 #

你的CDP调用语法完全符合规范，没有功能性错误，但有两个常见使用误区会额外增加开销：

• 不要在每个线程内单独创建、销毁CUDA流，流的创建和销毁本身就有不可忽略的设备端开销。如果一定要用动态并行，可以预先在主机端创建好流数组，作为参数传入父内核复用，能进一步降低少量开销，但依然无法抵消子内核启动的固有成本。
• 父内核用1个块200个线程的配置不是最优解，但这部分调度开销和子内核启动开销相比可以忽略，不会造成数量级的性能差。

3. 小粒度批量任务的优化建议 #

如果你的实际业务场景都是这类小尺寸矩阵批量处理，不要使用动态并行：

• 直接采用你最开始的基础实现即可，无需纠结少量线程闲置的问题，warp级的闲置成本远低于任何额外调度带来的开销。
• 如果矩阵尺寸差异极大，可以按尺寸对矩阵分组，给不同尺寸的矩阵组配置刚好匹配的线程数启动内核，进一步减少warp浪费，但实际收益通常不会超过10%。

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