[Bug已解决] 因在主机-host-代码中使用 CUDA_ARCH 宏导致 cudaErrorInvalidDeviceFunction 解决方案

[Bug已解决] 因在主机host代码中使用CUDA_ARCH宏导致 cudaErrorInvalidDeviceFunction 解决方案一、现象长什么样你在写自定义 CUDA 扩展.cu 文件里面用__CUDA_ARCH__这个宏做架构相关的分支但把它用在了主机host即非__global__/__device__代码里。运行时报cudaErrorInvalidDeviceFunction ... (某个 kernel 没找到对应架构的设备函数)即 pytorch/pytorch#177014。含义__CUDA_ARCH__是「设备代码编译期」才定义的宏表示正在为哪个 GPU 架构编译设备函数若你在 host 代码里用它的条件分支去决定加载/调用哪个 kernel会导致编译出的内核没有正确的设备函数运行时报cudaErrorInvalidDeviceFunction。 本文聚焦__CUDA_ARCH__的正确用法、为什么在 host 里用它会出问题、自定义 CUDA 扩展怎么正确地做架构分支。二、背景CUDA_ARCH是什么__CUDA_ARCH__是 NVCC 在编译设备代码device code时定义的宏值如750sm_75、860sm_86、900sm_90。它只在编译__device__/__global__函数体的时候存在用来写「针对不同架构走不同设备代码」的分支__global__ void my_kernel() { #if __CUDA_ARCH__ 700 // 针对 sm_70 的设备代码 #else // 老架构 #endif }关键点__CUDA_ARCH__只在「设备函数体内」有意义。在主机代码普通函数、全局作用域的 host 逻辑里__CUDA_ARCH__是未定义的或只反映「编译 host 侧时」的状态取决于分离编译设置。三、为什么在 host 代码用CUDA_ARCH会崩错误写法在 host 代码里用__CUDA_ARCH__决定「调用哪个 kernel 指针」或「是否启动某 kernel」// 错误在 host 代码里用 __CUDA_ARCH__ void launch() { #if __CUDA_ARCH__ 800 my_kernel_sm80...(); // host 侧 #if 分支 #else my_kernel_fallback...(); #endif }问题NVCC 编译时host 代码与 device 代码分两次编译。host 侧编译时__CUDA_ARCH__可能未定义 / 取默认值于是#if __CUDA_ARCH__ 800在 host 侧恒为假 → 只编译进 fallback 路径实际运行时如果你的 GPU 是 sm_86程序却因为 host 侧分支选了 fallback kernel而 fallback kernel 的设备函数没为 sm_86 编译或被分离编译丢弃→ 调用时找不到对应架构的设备函数 →cudaErrorInvalidDeviceFunction或者反过来host 侧认为该用 sm80 kernel但该 kernel 的设备代码没被正确包含 → 同样的「设备函数缺失」。 本质__CUDA_ARCH__是设备编译期的概念被误用在 host 编译期做运行时决策导致「实际 GPU 架构」与「被编译进的设备函数」不匹配。四、最小可运行复现编译期 C/CUDA 示意下面用极简 CUDA 示意「host 里误用CUDA_ARCH」的错误模式实际崩需编译运行// demo_cuda_arch_host.cu (示意错误不要照抄) #include cstdio __global__ void kernel_sm80() { /* sm_80 专用 */ } __global__ void kernel_fallback() { /* 老架构 */ } // 错误在 host 函数里用 __CUDA_ARCH__ 决定启动哪个 kernel void launch_wrong() { #if __CUDA_ARCH__ 800 // host 侧编译时此宏可能未定义 → 恒假 kernel_sm801, 1(); #else kernel_fallback1, 1(); #endif } // 正确host 侧用运行时查询设备属性而非编译期宏 #include cuda_runtime.h void launch_right() { int dev; cudaGetDevice(dev); cudaDeviceProp prop; cudaGetDeviceProperties(prop, dev); int major prop.major, minor prop.minor; if (major * 10 minor 80) { kernel_sm801, 1(); } else { kernel_fallback1, 1(); } } int main() { launch_right(); cudaDeviceSynchronize(); return 0; }要点host 侧做架构分支应该用「运行时cudaGetDeviceProperties」或「编译出全架构内核 运行时 dispatch」绝不能用__CUDA_ARCH__。五、解决方案一host 侧用运行时设备属性推荐正确的架构分支放在运行时host 侧查询真实设备#include cuda_runtime.h void launch() { cudaDeviceProp prop; int dev; cudaGetDevice(dev); cudaGetDeviceProperties(prop, dev); int capability prop.major * 10 prop.minor; if (capability 80) { kernel_sm80grid, block(); } else { kernel_fallbackgrid, block(); } }这样无论编译时怎么分离编译运行时根据实际 GPU 选对 kernel不会「设备函数缺失」。六、解决方案二编译全架构内核 运行时 dispatch更彻底用-gencode archcompute_XX,codesm_XX把每个架构的设备函数都编译进去运行时让 CUDA 驱动自动选匹配的# 编译命令为多个架构都生成代码 nvcc -gencode archcompute_70,codesm_70 \ -gencode archcompute_80,codesm_80 \ -gencode archcompute_90,codesm_90 \ my_ext.cu -o my_ext.oPyTorch 扩展里通过TORCH_CUDA_ARCH_LIST控制export TORCH_CUDA_ARCH_LIST7.0;8.0;9.0 # 编译多架构 python setup.py install这样每个 kernel 都针对 70/80/90 编译了设备函数运行时驱动自动 dispatch 到匹配架构无论 host 侧怎么 dispatch 都不会「设备函数缺失」。代价编译产物变大、编译更慢。七、解决方案三CUDA_ARCH只用在设备函数体内若你的分支本质是「设备端算法选择」同一 kernel 内不同架构走不同指令则把__CUDA_ARCH__留在__global__/__device__函数体内__global__ void my_kernel(float* x) { #if __CUDA_ARCH__ 800 // sm_80 的设备代码路径如用新指令 #else // 老架构设备代码 #endif }注意这种情况下NVCC 会为每个被编译的架构分别实例化这个 kernel 的设备代码每个架构取对应分支运行时驱动按架构选。这才是__CUDA_ARCH__的合法用途——设备函数体内而非 host 调度逻辑。八、解决方案四用 PyTorch 的架构抽象CUTLASS / 模板替代手写宏如果你用 PyTorch C 扩展优先用ATen 的at::cuda::getCurrentDeviceProperties()或模板分发而非裸__CUDA_ARCH__#include ATen/cuda/CUDAContext.h void launch() { int dev at::cuda::current_device(); const cudaDeviceProp* prop at::cuda::getDeviceProperties(dev); int capability prop-major * 10 prop-minor; if (capability 80) kernel_sm80...(); else kernel_fallback...(); }at::cuda::getDeviceProperties是 PyTorch 暴露的运行时设备查询比手写cudaGetDeviceProperties更贴合扩展环境。九、排查清单自定义 CUDA 扩展报cudaErrorInvalidDeviceFunction且用了__CUDA_ARCH__→ 确认 #177014 类误用。查是否在 host 代码普通函数/launch里用__CUDA_ARCH__做分支→ 这是根因。改运行时查询cudaGetDeviceProperties/at::cuda::getDeviceProperties决定 kernel方案一/四。编译多架构TORCH_CUDA_ARCH_LIST7.0;8.0;9.0让每个 kernel 都有匹配设备函数方案二。__CUDA_ARCH__只留在__global__/__device__函数体内做设备算法选择方案三。自检运行时打印capability确认 dispatch 的 kernel 真实存在该架构代码。十、小结cudaErrorInvalidDeviceFunction error due to use of __CUDA_ARCH__ in host code#177014的本质是__CUDA_ARCH__是 NVCC 在「编译设备代码」时定义的宏表示目标 GPU 架构只在__device__/__global__函数体内有意义。若把它用在 host 代码里做「启动哪个 kernel」的分支host 侧编译时宏未定义/取值错误导致实际 GPU 架构与「被编译进的设备函数」不匹配运行时找不到对应架构的设备函数 →cudaErrorInvalidDeviceFunction。 应对host 侧用运行时查询cudaGetDeviceProperties/at::cuda::getDeviceProperties按真实设备选 kernel绝不用__CUDA_ARCH__做 host 调度编译多架构TORCH_CUDA_ARCH_LIST7.0;8.0;9.0让每个 kernel 都有匹配的设备函数运行时自动 dispatch__CUDA_ARCH__仅限设备体内只用来在 kernel 内选设备算法不做 host 调度用 PyTorch 抽象优先at::cuda::getDeviceProperties而非裸 CUDA API。 记住__CUDA_ARCH__是「设备编译期」宏不是「运行时」属性。host 代码想知道 GPU 架构必须查cudaGetDeviceProperties而不是看编译期宏。