CUDA 合作组：链接错误答案

【问题标题】：CUDA Cooperative Groups : Linking errorCUDA 合作组：链接错误
【发布时间】：2020-04-16 06:12:29
【问题描述】：

在阅读了 CUDA 9 中的合作组之后，我一直在尝试在网格级别进行同步。

我正在使用 Visual Studio 2017、GTX 1060 和 CUDA 9.1。

我修改了我的代码如下：

__global__ void ExplicitKernel_American(/* ... */) {
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    auto grid = cooperative_groups::this_grid();
    if (i < sizeS) {
        //...
        for (int j = 1; j < sizeT; ++j) {
            // ...
            grid.sync(); // __syncthreads();
        }
    }
}

而且，正如文档中所述，我这样称呼我的内核：

void* Explicit_Args[] = { &PDE_Grid, /* ... */, &sizeS, &sizeT };
cudaLaunchCooperativeKernel(
    (void*)ExplicitKernel_American, 
    dim3((sizeS + TPB - 1) / TPB), 
    dim3(TPB),  
    Explicit_Args
); // TPB being 256...

不幸的是，我在内核中添加“网格”部分后立即出现链接错误。

error LNK2001: unresolved external symbol __fatbinwrap_38_cuda_device_runtime_compute_70_cpp1_ii_8b1a5d37
fatal error LNK1120: 1 unresolved externals

我设置了 -rdc=true 和 sm_61 但找不到它为什么不起作用...有什么想法吗？

非常感谢！

【问题讨论】：

标签： cuda linker-errors nvcc gpu-cooperative-groups

【解决方案1】：

使用协作内核启动（协作网格 - CG）需要 Pascal 或 Volta GPU，并且需要在 TCC 模式下运行的 Linux 或 Windows 设备。如果你测试设备属性结构中的deviceProp.cooperativeLaunch属性，我想你会发现你的GPU在WDDM模式下不支持它。

在尝试使用协作网格启动之前，最好在代码中测试此属性。

但是，您要问的问题是编译/链接问题。为此，我的建议是研究 CG（协作网格）示例代码，例如 6_Advanced/reductionMultiBlockCG。对于网格同步，绝对需要设置-rdc=true（即启用可重定位设备代码链接）。根据您设置-rdc=true 的方式，它可能无法正确应用于您的项目。概述了正确的方法here

这里的近端问题似乎是您没有正确链接到设备运行时库，例如-lcudadevrt

【讨论】：

您好，确实，问题是我缺少-lcudadevrt。有没有办法让这个功能发挥作用？我知道它应该可以在 Linux 平台上运行。另外，为什么存在这种限制？尤其是在 WINdows 上，它没有被记录在案！
实现它的方法我已经在我的回答中介绍过了。您目前无法使其在 WDDM 模型下工作。 WDDM 模型在 CUDA 和设备活动之间强加了一个操作系统管理的层，从而阻止了使此功能工作所需的控制级别。关于文档，这是当前文档中的疏忽，我希望在未来的 CUDA 版本中解决。