PTX 内核名称修改

Question

当从

ptx

文件编译内核时，我无法链接我的 Cuda 程序。

main.cu:

extern 
__global__ void kernel(int, float*);

int main()
{
    ...
    kernel<<<...>>>(...);
    ...
}

内核.cu

__global__
void kernel(int n, float* p)
{
    ...
}

如果我像下面这样编译，我没有问题，并且得到一个可执行文件：

nvcc -dc main.cu kernel.cu --gpu-architecture=sm_70
nvcc -dlink main.o kernel.o --gpu-architecture=sm_70 -o dlink.o
g++ dlink.o main.o kernel.o -lcudart

如果我像下面这样编译（通过生成

ptx

），我会收到错误：

nvcc -ptx kernel.cu --gpu-architecture=sm_70
nvcc -dc main.cu kernel.ptx --gpu-architecture=sm_70
nvcc -dlink main.o kernel.o --gpu-architecture=sm_70 -o dlink.o
g++ dlink.o main.o kernel.o -lcudart

错误：

main.o: In function `main':
tmpxft_0000b5ce_00000000-5_main.cudafe1.cpp:(.text+0x4789): undefined reference to `kernel(int, float*)'
tmpxft_0000b5ce_00000000-5_main.cudafe1.cpp:(.text+0x497e): undefined reference to `kernel(int, float*)'
collect2: error: ld returned 1 exit status

我正在遵循

CUDA_Compiler_Driver_NVCC.pdf

的示例。

我需要做什么来修复错误？

（这是 CUDA 10.2）。

Answer 1

如果您想编写自己的 PTX（或修改 PTX），则正确使用的 CUDA 方法是 CUDA 驱动程序 API 和相关的编译流程。

CUDA

vectorAddDrv

示例代码 拥有您需要的所有管道和工作流程。

Answer 2

不是对上述问题的直接回答。但是，当人们生成 ptx 文件时，目的是利用驱动程序 API 和 JIT（手动配置和启动内核）而不是运行时 API。至少我们在生产代码中是这样做的：

1. 对于全局内核，您可以对其进行标记

   extern "C"

   ，这样名称就不会被破坏
2. 当您调用内核时，您不会使用

   my_kernel<<<...>>>

   而是使用所有必要的参数来遍历

   cuLaunchKernel

第 1 条的意图是 s.t.在第 2 条中，您可以通过字符串（或字符串文字）中的名称告诉 JIT 编译器内核，例如 cuLaunchKernel 配置中的“my_kernel”。