torch.compile无法优化细粒度内存访问或kernel内部逻辑，而C++/CUDA扩展可控制寄存器、共享内存及warp同步；常见信号包括kernel launch过多、aten::小操作堆积、GPU利用率低但显存带宽满载。
为什么直接用 torch.compile 不够，必须写 C++ 扩展？
当你发现某个自定义算子在 Python 层反复调用 torch.tensor 操作、频繁进出 CUDA 上下文，或涉及细粒度内存访问（如 block-wise softmax、稀疏索引 gather），torch.compile 很难将其完全融合优化——它不触达 kernel 内部逻辑。此时 C++/CUDA 扩展是唯一能控制寄存器使用、共享内存布局、warp-level 同步的路径。
常见信号包括：nvprof 显示 kernel launch 过多、torch.autograd.profiler 报出大量 aten:: 小操作、GPU 利用率长期低于 30% 但显存带宽跑满。
用 torch.utils.cpp_extension.load 快速验证原型是否可行
别一上来就配 setup.py 或 CMakeLists.txt。先用动态加载绕过构建链路，5 分钟确认 C++ 逻辑是否跑通：
- 写一个
add_cuda.cu，含add_kernel和addhost 函数，确保调用AT_ASSERTM做 shape 检查 - Python 中用
load加载：from torch.utils.cpp_extension import load mod = load(name="add", sources=["add_cuda.cu"], extra_cuda_cflags=["-O3"])
- 注意：
extra_cuda_cflags必须显式加-O3，否则 PTX 编译不内联，性能差 3 倍以上 - 首次运行会触发
nvcc编译并缓存到~/.cache/torch_extensions，改代码后需清缓存或改name
PyTorch C++ API 的三个关键陷阱
不是所有 PyTorch C++ 头文件都能直接用；很多文档没写的隐含约束会导致 segfault 或梯度错乱：
立即学习“Python免费学习笔记（深入）”；
-
at::Tensor的.data_ptr()返回的是 raw pointer，但必须配合.device()和.dtype()手动校验——CUDA kernel 里用错float*当half*不报错，只输出 NaN - 自定义
backward时，不能直接返回tensor.clone()，得用tensor.detach().requires_grad_(True)，否则反向图断裂 - 若算子需支持
torch.compile，必须在 C++ 中显式注册torch::RegisterOperators().op("myop::forward", ...)，否则编译器无法 trace 到该 op
如何让 CUDA kernel 真正跑满 GPU？
写完 kernel 不等于高性能。观察 nsys profile 输出里的 Occupancy 和 Warp Execution Efficiency：
- 每个 block 启动线程数别硬写
256，用cudaOccupancyMaxPotentialBlockSize动态查询最优值 - 避免 global memory 随机访问：把索引数组预加载进 shared memory，再用
__syncthreads()对齐 warp - 混合精度要分两套 kernel：一套
half输入 +floataccumulator +half输出，另一套全float，不能靠__half2float在循环里强转 - kernel 启动参数中的 grid size 计算必须用
(n + block_size - 1) / block_size，整除写法n // block_size在 n 不整除时丢数据
最常被忽略的是 stream 同步：即使单个 kernel，也要传入 at::cuda::getCurrentCUDAStream() 而非默认 0 stream，否则和 PyTorch 主 stream 竞争导致隐式同步。
