dynamic并行 – 启动许多小内核是非常缓慢的

我正在尝试使用dynamic并行来改进我在CUDA中的algorithm。 在我原来的CUDA解决scheme中,每个线程计算每个块通用的编号。 我想要做的是首先启动一个粗略(或低分辨率)的内核,在这个线程中,线程只计算一次公共值(就像每个线程代表一个块一样)。 然后每个线程创build一个1个块(16×16线程)的小网格,并为其传递公共值启动一个子内核。 从理论上讲,它应该更快,因为它可以节省许多冗余操作。 但实际上,解决scheme的工作非常缓慢,我不知道为什么。

这是代码,非常简单,只是想法。

__global__ coarse_kernel( parameters ){ int common_val = compute_common_val(); dim3 dimblock(16, 16, 1); dim3 dimgrid(1, 1, 1); child_kernel <<< dimgrid, dimblock >>> (common_val, parameters); } __global__ child_kernel( int common_val, parameters ){ // use common value do_computations(common_val, parameters); } 

child_kernels的数量很多,每个线程一个,并且必须有大约400×400线程。 据我所知,GPU应该并行处理所有这些内核? 或者子内核按顺序处理? 我的结果显示,性能比我原来的解决scheme慢10倍以上。

Solutions Collecting From Web of "dynamic并行 – 启动许多小内核是非常缓慢的"

启动内核,无论是父母还是孩子都有成本。 如果您的子内核没有提取太多的并行性,而且对于非并行机制没有太大的好处,那么您的微弱优势可能会被子内核启动开销所抵消。

在公式中,让to成为执行子内核的开销,执行时间和执行相同代码的时间,而不需要动态并行的帮助。 使用动态并行的加速是ts/(to+te) 。 也许(但是这不能从你的代码中得到) te<ts但是te,ts<<to ,所以ts/(to+te)约为(ts/to)<1 ,你观察到一个减速而不是加速。