【问题标题】:Can I use cooperative_groups::sync(grid) in child kernel (CUDA dynamic parallelism)?我可以在子内核(CUDA 动态并行)中使用合作组::同步(网格)吗?
【发布时间】:2022-01-26 21:40:47
【问题描述】:
我有一个使用协作组的 reduce 内核(包含 cg::sync(grid)、cg::sync(cta)、grid.thread_rank() 等)。当从主机启动reduce内核时,它可以正常工作。当reduce内核从另一个内核(父内核)启动(作为子内核)时,我得到了未指定的启动失败错误。如果我从子内核中删除 cg::sync(grid) ,则没有错误。
所以,我的问题是:
我可以在子内核中使用cooperative_groups::sync(grid) 吗(CUDA 动态并行)?
【问题讨论】:
标签:
dynamic
parallel-processing
cuda
【解决方案1】:
到目前为止,CUDA 扩展“合作组”和 CUDA 扩展“动态并行”的网格同步功能不能混用。
引用cudaLaunchCooperativeKernel 上的CUDA 11.6.0 documentation,这是网格同步所需的:
内核无法使用 CUDA 动态并行。
编辑:此答案的早期版本指出,根本不能混合两个 CUDA 扩展。如果使用协作组而不使用网格同步功能,则不必使用cudaLaunchCooperativeKernel,因此应该能够使用动态并行。我更正了答案以反映这一点。