VS2019下CUDA Driver API动态并行代码报invalid PTX错误如何排查

开发环境说明 #

• 开发工具：VS2019
• 实现方式：C语言基于CUDA Driver API接口
• 错误现象：编译无报错，运行时加载模块时报invalid PTX错误

参考代码 #

主机端代码 #

CUdevice* gpu_initialize(CUdevice *dev_ptr, int int_of_dev);        //初始化函数签名

// 主机端代码
int main()
{
printf("Nested Parallelism : Hello World (Driver API)\n");

CUdevice device;
if (gpu_initialize(&device,0) == NULL) {                //初始化device[0]并填充设备指针内容
    printf("Error Initializing GPU... exiting");
    exit(-1);
}

CUresult result;
char* error_string;

// 创建上下文
CUcontext  context;
unsigned int flags = CU_CTX_SCHED_YIELD;            //设置为GPU等待CPU时主动让出的模式，会提高延迟
if ((result = cuCtxCreate(&context,flags,device)) != CUDA_SUCCESS) {
    if (cuGetErrorName(result, &error_string) != CUDA_ERROR_INVALID_VALUE) {
        printf("Error creating context : %s\n", error_string);
        return -1;
    }
    else {
        printf("Unknown error creating context\n");
        return -1;
    }
}

// 指定文件名加载模块
CUmodule module;
char filename[] = "C:\\Users\\gautam\\OneDrive\\Projects\\VS 2019\\repos\\learn_cuda_nesting\\dynamic_parallelism\\x64\\Debug\\nestedHelloWorld.cu.obj";        //VS2019编译产物路径

if ((result = cuModuleLoad(&module,filename)) != CUDA_SUCCESS) {
    if (cuGetErrorName(result, &error_string) != CUDA_ERROR_INVALID_VALUE) {
        printf("Error loading Module using filename %s: %s \n",filename,error_string);
        return -1;
    }
    else {
        printf("Unknown error loading Module from filename %s\n",filename);
        return -1;
    }
}
else printf("Successfully loaded module %s\n", filename);

// 从模块中加载函数
CUfunction function;
char        function_name[120];
strcpy_s(function_name,120,"nestedHelloWorld");
if ((result = cuModuleGetFunction(&function,module,function_name)) != CUDA_SUCCESS) {
    if (cuGetErrorName(result, &error_string) != CUDA_ERROR_INVALID_VALUE) {
        printf("Error loading function %s: %s\n",function_name, error_string);
        return -1;
    }
    else {
        printf("Unknown error loading function %s\n",function_name);
        return -1;
    }
}
else printf("Successfully loaded function %s\n", function_name);

// 设置核函数网格参数

int size = 8;
int blocksize = 8;   // 初始块大小
int igrid = 1;
int threads_per_block_x = 8;
int threads_per_block_y = 1;
int blocks_per_grid_x = (size + threads_per_block_x - 1) / threads_per_block_x;
int blocks_per_grid_y = 1;

// 启动第一个核函数
// 场景1：以8x1的网格大小调用核函数
void* args[] = {(void *)&size, (void *)&igrid};

if ((result = cuLaunchKernel(function, blocks_per_grid_x, blocks_per_grid_y, 1, threads_per_block_x, threads_per_block_y, 1, 0, 0, args, 0)) != CUDA_SUCCESS) {
    if (cuGetErrorName(result, &error_string) != CUDA_ERROR_INVALID_VALUE) {
        printf("Error launching kernel %s: %s\n",function_name, error_string);
        result = cuCtxDestroy(context);
        return -1;
    }
    else {
        printf("Unknown error launching kernel %s\n", function_name);
        result = cuCtxDestroy(context);
        return -1;
    }
}
else printf("CUDA kernel launch with (%d,%d) blocks per grid, each with (%d,%d) threads per block\n", blocks_per_grid_x, blocks_per_grid_y, threads_per_block_x, threads_per_block_y);

result = cuCtxSynchronize();        //等待核函数执行完成
result = cuCtxDestroy(context);
exit(0);
}

设备端.cu代码 #

extern "C" __global__ void nestedHelloWorld(int const iSize, int iDepth)
{
int tid = threadIdx.x;
printf("Recursion=%d: Hello World from thread %d block %d\n", iDepth, tid,blockIdx.x);

// 递归终止条件
if (iSize == 1) return;

// 块大小减半
int nthreads = iSize >> 1;

// 0号线程递归启动子网格
if(tid == 0 && nthreads > 0)
{
    nestedHelloWorld<<<1, nthreads>>>(nthreads, ++iDepth);
    printf("-------> nested execution depth: %d\n", iDepth);
}
}

错误原因与修复方案 #

1. 加载的文件类型错误 #

你当前用cuModuleLoad加载的是VS编译生成的.obj目标文件，该文件是给Runtime API主机端链接使用的混合格式文件，不是Driver API支持的纯设备镜像文件（cubin/PTX），所以会识别为无效PTX。
修复方法：

• 打开VS项目属性，进入CUDA C/C++ -> Common，勾选生成GPU设备代码文件，或者手动给nvcc添加编译参数-cubin，编译后会生成独立的.cubin设备镜像文件
• 修改代码中filename路径，指向生成的.cubin文件

2. 动态并行编译配置缺失 #

动态并行要求设备代码必须开启可重定位设备代码编译，且完成设备端链接：

• 确认CUDA C/C++ -> 主机 -> 生成可重定位设备代码选项设置为是（-rdc=true）
• 确认链接器附加依赖项中已经添加cudadevrt.lib，且库路径包含对应CUDA版本的x64库目录
• 确认CUDA 链接器 -> 常规 -> 执行设备链接选项设置为是

3. 计算架构不匹配 #

你当前配置的目标架构是compute_75, sm_75，仅支持图灵架构及以上显卡（GTX16xx、RTX20xx及更新系列），如果你的显卡实际算力低于75，生成的设备代码无法在当前设备加载，也会报invalid PTX错误。
修复方法：查询你的显卡对应算力，修改编译选项中的目标架构为匹配的版本，可在GPU初始化逻辑中添加代码打印当前设备的算力版本用于校验。

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