我可以在子内核中使用cooperative_groups::sync(grid)（CUDA动态并行）吗？

Question

我有一个使用合作组的reduce内核（包含

cg::sync(grid)

，

cg::sync(cta)

，

grid.thread_rank()

等）。当reduce内核从主机启动时，它可以正常工作。当reduce内核从另一个内核（父内核）启动（作为子内核）时，我收到“未指定的启动失败”错误。如果我从子内核中删除

cg::sync(grid)

，就不会出现错误。

我可以在子内核中使用

cooperative_groups::sync(grid)

（CUDA动态并行）吗？

Answer 1

迄今为止，CUDA 扩展“合作组”和 CUDA 扩展“动态并行”的网格同步功能无法混合。

引用cudaLaunchCooperativeKernel上的CUDA 11.6.0

文档

，这是网格同步所需要的：

内核无法利用 CUDA 动态并行性。

编辑：这个答案的早期版本指出，根本不能混合使用两个 CUDA 扩展。当使用协作组而不使用网格同步功能时，不必使用

cudaLaunchCooperativeKernel

，因此应该能够使用动态并行性。我更正了答案以反映这一点。