CUDA内核中禁用std::string等标准容器、异常处理、RTTI、虚函数、非空构造/析构的自定义类型、可变参数模板（实参不满足约束时）、__CUDA_ARCH__用于签名控制、内存空间修饰符用于类成员或形参，以及未标注__device__的函数调用。
CUDA内核中禁用哪些C++特性？答案很直接：不是“哪些能用”，而是“哪些明确不能用”——尤其在 __global__ 和 __device__ 函数里，很多看似无害的 C++ 特性会静默失效、编译报错，或引发运行时未定义行为。
哪些 C++ 特性在 __device__ 代码里根本不可用？
官方明确禁止（非建议，是硬性限制）：
-
std::string、std::vector、std::map等标准容器 —— 设备端无堆管理，且不支持动态内存分配； - 异常处理（
try/catch）、RTTI（dynamic_cast、typeid）—— SM 指令集不支持，编译直接失败； - 虚函数调用（含
virtual、override）—— 设备端无 vtable 支持，即使类定义在主机端也禁止调用； - 非空构造/析构函数的自定义类型 —— 若
class或struct带有用户定义的构造函数、析构函数或拷贝操作符，就不能作为__device__变量成员或 kernel 参数； - 可变参数模板（
template<typename... Args>）在设备端实例化受限 —— 仅当所有实参类型本身满足设备端约束时才允许，否则链接时报undefined reference to __nv_...；
__CUDA_ARCH__ 宏误用：看似条件编译，实则埋雷
开发者常想用它做 host/device 分支，但这是高危操作：
- 不能用于决定变量/函数签名 —— 如
typedef、函数形参、模板参数推导等，会导致 host/device 编译结果不一致； - 不能用于
__global__函数模板的显式实例化 —— 下面这段代码在 CUDA 12.5 中仍报错：template<typename T> __global__ void kern(T x) { /* ... */ } #if !defined(__CUDA_ARCH__) kern<int><<<1,1>>>(42); // ❌ 错误：host 端调用 device-only 模板实例 #endif - 头文件中若用
__CUDA_ARCH__控制函数实现，多个 .cu 文件分别编译后链接，可能因 arch 不同导致符号冲突；
内存空间描述符的误用场景
__device__、__shared__、__constant__、__managed__ 这四类修饰符有严格使用边界：
立即学习“C++免费学习笔记（深入）”；
- 不能修饰 class 成员 —— 即使该 class 仅用于 device，其成员也不能加
__device__； - 不能出现在函数形参列表中 ——
void foo(__device__ int* p)是非法语法； - 在 device 函数中，除
__shared__外，其他三类变量必须声明为static或extern（单独编译模式下），否则 NVCC 报error: variable with memory space qualifier must be static or extern； -
__shared__变量若为类类型，该类必须满足：无构造函数、无析构函数、无虚函数、所有成员为 POD 类型；
CUDA 12.5 的“松动”与陷阱
虽然 CUDA 12.5 开始支持设备端 lambda（[] __device__ (int x) { return x * 2; }），但要注意：
- lambda 必须显式标注
__device__，漏写即默认 host-only，运行时崩溃； - 捕获列表仅支持值捕获（
[x]），不支持引用捕获（[&x]）或 this 捕获； - 不能嵌套定义 lambda，也不能在 lambda 内递归调用自身；
- 若 lambda 捕获了 host 端变量（如
int a = 5; auto f = [a] __device__ { return a; };），该变量会在 kernel 启动时按值复制进 device 寄存器 —— 无法反映 host 端后续修改；
真正容易被忽略的，是那些“编译通过但行为错乱”的情况：比如在 __device__ 函数里调用一个没加 __device__ 标记的普通函数，NVCC 可能静默降级为 host 调用（取决于编译模式），结果 kernel 执行到那行就卡死或返回错误码 cudaErrorInvalidDeviceFunction。这类问题不会在编译期暴露，得靠 cudaGetLastError() 和 Nsight Compute 逐行排查。
