c++ - 在 CUDA 内核中调用内核

Question

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

__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 时它给出的错误:

Error MSB3721 The command ""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" exited with code 1.

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

Answer 1

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

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

2. 在生成的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];
    }
   

3. 在 Project...Properties...CUDA C++...Common 中将 Generate Relocatable Device Code 设置为“Yes”

4. 在 Project...Properties...CUDA Linker...General 中将 cudadevrt.lib 添加到 Additional Dependencies

5. 构建或重建项目，然后您应该会看到如下输出:

    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 版本而异，但您可以通过搜索找到说明，例如 this .目的是增加详细程度，以便 VS 会在出现错误时向您显示 nvcc 生成的实际输出。

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

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