因为torch.cuda.is_available()仅检测CUDA运行时可用性，不保证nvcc编译器、头文件、链接库等开发工具链存在，而CUDA扩展编译必须依赖完整CUDA Toolkit。
为什么不能直接用 torch.cuda.is_available() 就认为能写 CUDA 扩展
很多开发者看到 PyTorch 支持 CUDA，就默认本地环境可以直接编译 CUDA 算子，结果在 python setup.py install 阶段报 NVCC not found 或 no CUDA toolset found。根本原因是：PyTorch 的 C++ 扩展（torch.utils.cpp_extension）依赖系统级 CUDA 工具链，而非仅靠 PyTorch 自带的 CUDA 运行时。
必须确认三件事：
-
nvidia-smi能查到驱动版本，且 ≥ 对应 CUDA 版本要求（例如 CUDA 12.1 要求驱动 ≥ 530） -
nvcc --version可执行，且版本与 PyTorch 编译时使用的 CUDA 版本一致（可通过torch.version.cuda查） - Windows 用户额外需要匹配的 Visual Studio 工具集（如 CUDA 12.x 要求 VS 2019/2022，且需安装 CMake Tools 和 Windows SDK）
如何用 CUDAExtension 正确组织文件结构
常见错误是把 .cu 文件和 Python 脚本混放，或忽略头文件路径。PyTorch 要求显式声明 CUDA 源、头文件、include 路径和链接库，否则编译器找不到 AT_ASSERT 或 cudaStream_t。
标准最小结构如下：
立即学习“Python免费学习笔记（深入）”；
myop/ ├── setup.py ├── myop.cpp // C++ binding（调用 torch::RegisterOperators） ├── myop_cuda.cu // CUDA kernel + launch wrapper └── myop_cuda.h // 声明 kernel 函数和 launch 函数（extern "C"）
关键点：
-
setup.py中必须用CUDAExtension（不是CppExtension），并传入include_dirs（含torch/include和torch/include/torch/csrc/api/include） -
myop_cuda.h里所有 CUDA 函数声明前加extern "C"，避免 C++ name mangling 导致 Python 无法 dlsym -
myop.cpp中用torch::RegisterOperators().op("myop::forward", &myop_forward)注册，函数签名必须严格匹配（如输入为const torch::Tensor&，输出为torch::Tensor）
torch::Tensor 如何安全地映射到 CUDA device pointer
新手常直接对 tensor.data_ptr<float>() 做指针运算，却忽略 tensor 是否 contiguous、是否在 GPU 上、dtype 是否匹配——这会导致非法内存访问或静默错误结果。
正确做法分三步：
- 强制 contiguous：
input.contiguous()，再取指针；否则用input.unsafe_chunk()等高级 API（不推荐） - 检查设备：
input.is_cuda()必须为 True，否则抛std::runtime_error("expected CUDA tensor") - 类型断言：
AT_ASSERTM(input.scalar_type() == at::ScalarType::Float, "only float supported")（AT_ASSERTM是 PyTorch 提供的安全宏）
示例片段（在 myop_cuda.cu 的 launch 函数中）：
const float* input_ptr = input.data_ptr<float>(); const int64_t n = input.numel(); myop_kernel<<<(n + 255) / 256, 256, 0, stream>>>(input_ptr, output_ptr, n);
为什么 torch.compile 无法加速自定义 CUDA 算子
PyTorch 2.0+ 的 torch.compile 默认只优化 TorchScript IR 中的原生算子（如 aten::add, aten::matmul），不会重写或融合你注册的 myop::forward。它看到的是一个 opaque operator call，直接 bypass 优化流程。
若想获得图优化收益，目前只有两条路：
- 将逻辑尽量用原生 PyTorch 算子组合（如用
torch.nn.functional.embedding替代手写 embedding lookup） - 手动实现
torch._dynamo.backends.register_backend并提供自己的 lowering 规则（极少见，文档极少，需深入理解 Inductor 的 FX Graph）
更现实的做法是：专注 kernel 本身的优化（shared memory 使用、bank conflict 避免、warp-level primtives），而不是依赖 torch.compile。
最易被忽略的一点：CUDA kernel 的 launch 参数（grid/block）若写死，会在不同 batch size 下严重低效；务必根据 tensor.numel() 动态计算，别抄网上静态示例。
