如何检查CUDA内核支持的计算能力及编译架构匹配异常排查

核心问题 #

• 是否可以在运行时检查当前程序编译所针对的CUDA计算能力？
• arch=compute_xx,code=sm_xx编译选项是否会设置可直接检查的宏定义？

背景 #

无法确保部署的二进制程序用户环境配置正确，希望在执行耗时初始化前，通过cudaGetDeviceProperties获取设备计算能力（CC），并与程序编译时的计算能力对比。目前只能通过尝试启动内核，再用cudaPeekAtLastError捕获“no kernel image is available for execution on the device”错误来退出。

补充问题（代码与异常） #

尝试用cudaFuncGetAttributes编写代码（main.cu）：

#include <iostream>

__global__ void cudaKernel()
{
    ;
}

int main()
{
    cudaFuncAttributes attr;
    cudaError_t err = cudaFuncGetAttributes(&attr, cudaKernel);
    if (err != cudaSuccess) {
        std::cout << "CUDA Error: " << cudaGetErrorString(err);
        return EXIT_FAILURE;
    }
    
    return EXIT_SUCCESS;
}

• 该代码在CUDA 11.2、GTX 1080（CC 6.1）的Linux机器上运行正常。
• 在CUDA 11.8、GT730（CC 3.5）的Windows机器上返回CUDA Error: invalid device function，官方文档说明该错误表示“请求的设备函数不存在或未针对正确的设备架构编译”。

用cuobjdump查看可执行文件，结果显示：

Fatbin ptx code:
================
arch = sm_35
code version = [7,8]
host = windows
compile_size = 64bit
compressed

Fatbin elf code:
================
arch = sm_35
code version = [1,7]
host = windows
compile_size = 64bit

困惑点：cuobjdump识别出架构为sm_35，但cudaFuncGetAttributes调用失败，而设备CC为3.5本应支持该架构，当前驱动版本456.71，官方兼容性文档显示支持CUDA 11.x。

解答 #

1. 编译选项对应的宏定义 #

arch=compute_xx,code=sm_xx编译选项本身不会自动生成可直接检查的宏，但你可以手动在编译时添加自定义宏，比如：

nvcc -arch=compute_35 -code=sm_35 -DCOMPUTE_CAP_MAJOR=3 -DCOMPUTE_CAP_MINOR=5 main.cu

之后在代码中就能直接使用COMPUTE_CAP_MAJOR和COMPUTE_CAP_MINOR这两个宏，获取编译时指定的计算能力版本，进而和cudaGetDeviceProperties拿到的设备CC做对比。

2. cudaFuncGetAttributes调用失败的原因 #

CUDA 11.8在Windows平台要求的最低驱动版本是472.50，你使用的456.71属于CUDA 11.0-11.1对应的驱动版本，无法完全兼容CUDA 11.8编译出的sm_35代码，导致驱动无法正确识别二进制中的设备函数，从而返回invalid device function错误。

3. 运行时检查编译目标CC的替代方案 #

• 手动定义宏（推荐）：前面提到的自定义宏方式简单直接，能快速实现编译CC和设备CC的对比。
• 解析Fatbin结构：通过cudaGetFatBinary获取程序的fatbin数据，手动解析其中的架构信息，但需要了解CUDA fatbin的格式，实现复杂度较高。
• 保留内核启动错误捕获：这是最可靠的通用方案，即使通过宏或解析拿到了编译目标CC，实际运行时驱动与二进制的兼容性仍可能出问题，启动内核并检查错误的方式能直接反映程序是否可正常运行。

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