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 或梯度错乱:
-
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 竞争导致隐式同步。
Python免费学习笔记(深入):立即使用
在学习笔记中,你将探索 Python 的核心概念和高级技巧!











