为什么不允许 cudaLaunchCooperativeKernel() 返回？答案

【问题标题】：Why is cudaLaunchCooperativeKernel() returning not permitted?为什么不允许 cudaLaunchCooperativeKernel() 返回？
【发布时间】：2021-04-01 00:07:05
【问题描述】：

所以我使用的是 GTX 1050，计算能力为 6.1 和 CUDA 11.0。我需要在我的程序中使用网格同步，所以需要cudaLaunchCooperativeKernel()。我检查了我的设备查询，因此 GPU 确实支持合作组。我无法执行以下功能

 extern "C" __global__ void test(int x) {
    if (x) {
       printf("%d", x);
       if (threadIdx.x == 0)
          test<<<1, 1>>>(--x);
    }
}

调用后，

cudaLaunchCooperativeKernel((void *)test, 1, 1, (void **) (&x));

收到错误“不允许操作”（代码为 800）。现在，当设备不支持协作组时返回（在这种情况下不支持）。那么，什么会导致这个问题呢？

【问题讨论】：

标签： cuda dynamic-parallelism gpu-cooperative-groups

【解决方案1】：

您的内核利用了动态并行性。但是，在通过cudaLaunchCooperativeKernel 启动的内核中不允许使用动态并行

这在运行时 API 的文档中有所提及。 https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__EXECUTION.html

【讨论】：

我明白了！那么在我的情况下，网格同步有什么解决方法吗？
我认为网格同步的唯一解决方法是拥有两个内核，一个用于同步前的部分，一个用于同步后的部分。