在CUDA内核中调用内核

Question

我正在尝试做这样的事情：

__global__ void foo()
{
    // do stuff
}

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

但是我遇到了错误 "kernel launch from __device__ or __global__ functions requires separate compilation mode"。

我尝试通过谷歌搜索来找到答案，看到一些结果提到了 "dynamic-parallelism"，它需要计算能力为 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 时它是如何工作的。

- paferllume

---

你能否提供你所看到的编译错误？ - cwharris

@cwharris 除了我已经写的内容之外，它只是说nvcc以代码1退出。 - paferllume

学习任何使用CUDA动态并行性的CUDA示例代码及其相关项目文件，例如cdpSimplePrint。 - Robert Crovella

@RobertCrovella 看起来我做对了，我继续调查发现问题似乎出在链接文件上，但我不完全明白在使用rdc时它是如何工作的。 - paferllume

Answer 1

使用MS VS 2019和CUDA 11.0，以下步骤可用于创建动态并行性（CDP）示例：

1. Create a new CUDA Runtime project

2. In the kernel.cu file that is generated, modify the kernel like so:

    __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];
    }
   

3. In Project...Properties...CUDA C++...Common set Generate Relocatable Device Code to "Yes"

4. In Project...Properties...CUDA Linker...General add cudadevrt.lib to Additional Dependencies

5. Build or rebuild the project, you should then see output like this:

    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 ==========
   

注意:

1. CUDA 11.0（及更高版本）仅针对支持CDP的设备。对于早期版本，您可能需要将设备代码生成目标设置为与支持CDP的GPU匹配（例如，compute_35，sm_35 ）

2. 在MS VS中，MSB3721错误本身并不是很有用。它只是表示“出了些问题”。要从Visual Studio中获取更有用的信息，您应该增加控制台输出的详细程度。确切的方法会因VS版本而异，但您可以通过搜索（例如此处）找到说明。目标是增加详细程度，以便在有错误时，VS会向您显示nvcc生成的实际输出。

3. 对于CUDA 11.0 / VS2019，添加cudadevrt.lib不是必需的，因为它已包含在项目中。但对于其他/较旧版本，则可能需要此项。

如果您仍然遇到问题，建议您增加详细程度以更好地了解确切的问题。您还应尝试按以上列出的步骤进行操作，确保您理解它们（即从新项目开始）。如果您仍遇到问题，请发布新问题，包括您的实际代码以及增加详细程度后的控制台编译输出。

- Robert Crovella

Answer 2

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

- paferllume