CUDA的PTX JIT编译器--entry参数输入规则及名称修饰含义问询

1. --entry参数接受什么输入？ #

该参数接受逗号分隔的一个或多个函数名称字符串，但这些名称必须是PTX代码中实际存在的、经过名称修饰后的函数标识。

2. 如何“以名称修饰后的形式指定”？ #

名称修饰是编译器为了区分不同上下文、不同签名的函数而生成的唯一标识（比如C++重载、模板实例化、CUDA核函数的参数差异都会触发修饰）。要获取并使用修饰后的名称，你可以：

• 用nvcc编译CUDA代码时添加--keep参数，保留生成的PTX文件，直接查看文件中.entry指令后的函数名（这就是修饰后的版本）。
• 用nvdisasm反编译CUDA二进制文件（如.cubin），从中提取目标函数的修饰名。
• 如果是动态生成的PTX代码，要确保你生成的PTX中.entry定义的名称，和传给--entry的名称完全一致。

举个实际例子：
假设你有如下CUDA代码：

__global__ void my_kernel(int a) { /* ... */ }

用nvcc --keep my_kernel.cu编译后，打开生成的my_kernel.ptx，会看到类似这样的定义：

.entry _Z10my_kerneli(
    .param .b32 _Z10my_kerneli_param_0
)
{
    /* PTX指令 */
}

这里的_Z10my_kerneli就是修饰后的函数名，JIT编译时就要把这个字符串传给--entry参数，比如用NVRTC API时：

const char* opts[] = {"--entry", "_Z10my_kerneli"};
nvrtcCompileProgram(prog, 2, opts);

3. 是否必须使用经过名称修饰的函数名？ #

是的，必须使用修饰后的名称。因为PTX代码中实际存在的函数标识就是修饰后的版本，JIT编译器只会匹配PTX里.entry定义的名称。如果传入未修饰的原名（比如my_kernel），编译器会找不到对应的入口函数，直接报错编译失败。

内容的提问来源于stack exchange，提问作者einpoklum