2

我正在尝试使用动态并行性来改进我在 CUDA 中的算法。在我最初的 CUDA 解决方案中，每个线程都会计算每个块通用的数字。我想要做的是首先启动一个粗略（或低分辨率）内核，其中线程只计算一次公共值（就像每个线程代表一个块一样）。然后每个线程创建一个由 1 个块组成的小网格（16x16 线程），并为它启动一个子内核以传递公共值。理论上它应该更快，因为它可以节省许多冗余操作。但在实践中，该解决方案的工作速度非常慢，我不知道为什么。

这是代码，非常简化，只是想法。

__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 的数量很多，每个线程一个，并且必须有大约 400x400 个线程。据我了解，GPU 应该并行处理所有这些内核，对吧？

或者子内核以某种方式按顺序处理？

我的结果表明，性能比我原来的解决方案慢 10 倍以上。

4

1 回答 1

3

启动内核是有成本的，无论是父内核还是子内核。如果您的子内核没有提取太多并行性，并且相对于它们的非并行对应物没有太多好处，那么您的微弱好处可能会被子内核启动开销所抵消。

在公式中，让to执行子内核的开销、te它的执行时间以及ts在没有动态并行的帮助下执行相同代码的时间。使用动态并行产生的加速是ts/(to+te). 也许（但这不能从您的代码中看出）te<ts但是te,ts<<to，这ts/(to+te)大约是(ts/to)<1您观察到减速而不是加速。

于 2014-01-08T22:10:23.800 回答