在 CUDA 内核中调用内核答案

【问题标题】：call kernel inside CUDA kernel在 CUDA 内核中调用内核
【发布时间】：2021-02-07 11:21:51
【问题描述】：

我正在尝试做类似的事情：

__global__ void foo()
{
    // do stuff
}

__global__ void boo()
{
    foo<<<m, n>>>();
}

但我收到错误“从 __device__ 或 __global__ 函数启动内核需要单独的编译模式”

我尝试在谷歌上搜索答案，我看到一些关于“动态并行性”的结果，它说它需要我拥有的计算能力 3 或更高（GTX 750 Ti 计算能力 5）。
我还需要打开“rdc”标志，虽然它确实使错误消失，但无论如何都会导致编译失败（即使我评论了所有内容）

那么我该如何实现我的目标或可能存在什么问题？
（使用 cuda 11.0）
我还添加了“cudadevrt.lib;cudart.lib;”在项目属性的链接器中输入

编辑：
rdc设置为true时给出的错误：

错误 MSB3721 命令 ""C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.0\bin\nvcc.exe" -dlink -o "x64\Debug\crimson cuda.device-link.obj" -Xcompiler "/EHsc /W3 /nologo /Od /Zi /Fdx64\Debug\vc142.pdb /RTC1 /MDd " -L"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.0\bin/crt" -L"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.0\lib\x64" cudadevrt.lib cudart.lib cudart_static.lib kernel32.lib user32.lib gdi32.lib winspool.lib comdlg32.lib advapi32. lib shell32.lib ole32.lib oleaut32.lib uuid.lib odbc32.lib odbccp32.lib -gencode=arch=compute_50,code=sm_50 -G --machine 64 x64\Debug\CrimsonNet.cu.obj x64\Debug\kernel。 cu.obj" 以代码 1 退出。

编辑 2：我继续调查，似乎是在链接文件时出现问题，我不完全理解使用 rdc 时它是如何工作的。

【问题讨论】：

你能发布你看到的编译错误吗？
@cwharris 除了我已经写过的内容外，它只是说 nvcc 以代码 1 退出
研究任何使用 CUDA 动态并行的 CUDA 示例代码及其关联的项目文件，例如 cdpSimplePrint
@RobertCrovella 看来我做得对，我继续调查，似乎问题出在链接文件时，我不完全理解使用 rdc 时它是如何工作的

标签： c++ cuda visual-studio-2019

【解决方案1】：

使用 MS VS 2019 和 CUDA 11.0，以下步骤允许我创建动态并行 (CDP) 示例：

创建一个新的 CUDA 运行时项目

在生成的kernel.cu文件中，修改内核如下：

 __global__ void child_kernel() {printf("hello\n");}

 __global__ void addKernel(int *c, const int *a, const int *b)
 {
     child_kernel << <1, 1 >> > ();
     int i = threadIdx.x;
     c[i] = a[i] + b[i];
 }

In Project...Properties...CUDA C++...Common set Generate Relocatable Device Code为“Yes”
在 Project...Properties...CUDA Linker...General 添加 cudadevrt.lib 到 Additional Dependencies

构建或重新构建项目，您应该会看到如下输出：

 1>------ Rebuild All started: Project: test23, Configuration: Debug x64 ------
 1>Compiling CUDA source file kernel.cu...
 1>
 1>C:\Users\Robert Crovella\source\repos\test23>"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.0\bin\nvcc.exe" -gencode=arch=compute_52,code=\"sm_52,compute_52\" --use-local-env -ccbin "C:\Program Files (x86)\Microsoft Visual Studio\2019\Professional\VC\Tools\MSVC\14.20.27508\bin\HostX86\x64" -x cu -rdc=true  -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.0\include" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.0\include"  -G   --keep-dir x64\Debug -maxrregcount=0  --machine 64 --compile -cudart static  -g   -DWIN32 -DWIN64 -D_DEBUG -D_CONSOLE -D_MBCS -Xcompiler "/EHsc /W3 /nologo /Od /Fdx64\Debug\vc142.pdb /FS /Zi /RTC1 /MDd " -o x64\Debug\kernel.cu.obj "C:\Users\Robert Crovella\source\repos\test23\kernel.cu"
 1>kernel.cu
 1>
 1>C:\Users\Robert Crovella\source\repos\test23>"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.0\bin\nvcc.exe" -dlink -o x64\Debug\test23.device-link.obj -Xcompiler "/EHsc /W3 /nologo /Od /Zi /Fdx64\Debug\vc142.pdb /RTC1 /MDd " -L"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.0\bin/crt" -L"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.0\lib\x64" cudadevrt.lib cudart_static.lib kernel32.lib user32.lib gdi32.lib winspool.lib comdlg32.lib advapi32.lib shell32.lib ole32.lib oleaut32.lib uuid.lib odbc32.lib odbccp32.lib cudart.lib  -gencode=arch=compute_52,code=sm_52 -G --machine 64 x64\Debug\kernel.cu.obj
 1>cudadevrt.lib
 1>cudart_static.lib
 1>kernel32.lib
 1>user32.lib
 1>gdi32.lib
 1>winspool.lib
 1>comdlg32.lib
 1>advapi32.lib
 1>shell32.lib
 1>ole32.lib
 1>oleaut32.lib
 1>uuid.lib
 1>odbc32.lib
 1>odbccp32.lib
 1>cudart.lib
 1>kernel.cu.obj
 1>   Creating library C:\Users\Robert Crovella\source\repos\test23\x64\Debug\test23.lib and object C:\Users\Robert Crovella\source\repos\test23\x64\Debug\test23.exp
 1>test23.vcxproj -> C:\Users\Robert Crovella\source\repos\test23\x64\Debug\test23.exe
 ========== Rebuild All: 1 succeeded, 0 failed, 0 skipped ==========

注意事项：

CUDA 11.0（及更高版本）仅针对支持 CDP 的设备。对于早期版本，您可能需要设置设备代码生成目标以匹配支持 CDP 的 GPU（例如 compute_35,sm_35）
在 MS VS 中，MSB3721 错误本身并没有那么有用。它只是表示“出了点问题”。要从 Visual Studio 获得更多有用的信息，您应该增加控制台输出的详细程度。执行此操作的确切方法因 VS 版本而异，但您可以通过搜索找到说明，例如 this。目的是增加详细程度，以便 VS 在出现错误时向您显示nvcc 生成的实际输出。
对于 CUDA 11.0/VS2019，不需要添加 cudadevrt.lib，因为它已经包含在项目中。对于其他/旧版本，可能是必要的。

如果您仍然遇到问题，我建议您增加详细程度以更好地了解确切的问题。您还应该完全尝试上面列出的步骤，以确保您理解它们（即从一个新项目开始）。如果您仍然遇到问题，请在您增加详细程度后使用您的实际代码以及控制台编译输出发布一个新问题。

【讨论】：

【解决方案2】：

我仍然不知道是什么导致了问题，但是在删除除驱动程序之外的所有与 Nvidia 相关的内容后，然后通过 CUDA 安装程序重新安装所有内容后，错误消失了，现在可以正常工作了。

【讨论】：