4. 核函数启动配置：<<>>语法，动态并行，启动参数优化
说到CUDA编程，最基础也最容易被忽视的就是核函数启动配置。说白了就是那对尖括号里的参数——<<<grid, block>>>。我见过太多人随便填个数字就跑，结果性能惨不忍睹。今天咱们好好聊聊这里面的门道。
4.1 启动配置的基本语法
先看最基础的写法：
kernel_name<<<gridDim, blockDim, sharedMem, stream>>>(args);四个参数，但很多人只用前两个。我个人习惯把四个都写上，哪怕sharedMem是0，stream是默认流。为什么？代码可读性更强，后期维护也方便。
| 参数 | 含义 | 典型范围 |
|---|---|---|
| gridDim | 网格维度（block数量） | 1 ~ 数百万 |
| blockDim | 线程块维度（线程数量） | 32 ~ 1024 |
| sharedMem | 动态共享内存大小（字节） | 0 ~ 48KB |
| stream | CUDA流 | 0 或 cudaStream_t |
核心原则：blockDim必须是32的倍数。这不是建议，是硬性要求。因为warp是32个线程一组，不是32的倍数会导致最后一个warp浪费。
4.2 block大小的选择艺术
block大小怎么选？我直接给结论：128或256是最稳妥的选择。
为什么？你想想看：
- block太小（比如32）：SM利用率低，调度开销大
- block太大（比如1024）：寄存器压力大，occupancy反而下降
我在项目中遇到过这样一个案例：一个矩阵乘法核函数，用256线程的block比用512线程快了15%。当时团队里有人不信，觉得线程越多并行度越高。其实不是的，线程多了每个线程分到的寄存器就少了，编译器不得不把变量 spill 到local memory，性能就崩了。
经验公式：block大小 = (寄存器使用量 × 线程数) ≤ 65536。用nvcc的--ptxas-options=-v可以查看寄存器使用情况。
4.3 grid大小的计算
grid大小说白了就是你要启动多少个block。常见做法：
int blockSize = 256;
int gridSize = (N + blockSize - 1) / blockSize;
kernel<<<gridSize, blockSize>>>(d_data, N);嗯，这里要注意：不要启动太多block。GPU能同时运行的block数量是有限的，取决于SM数量和每个SM能容纳的block数。启动几百万个block虽然语法上允许，但调度开销会让你怀疑人生。
我曾经犯过一个错误：处理1亿个元素，每个block只处理1个元素，启动了1亿个block。结果呢？比CPU还慢。后来改成每个block处理256个元素，性能提升了100倍。
4.4 动态并行：核函数里再调核函数
动态并行是CUDA里一个很酷的特性——在GPU上直接启动新的核函数。语法和host端一样：
__global__ void parent_kernel() {
if (threadIdx.x == 0) {
child_kernel<<<grid, block>>>();
}
}警告：动态并行有坑！子核函数启动是异步的，但父核函数不会自动等待子核函数完成。你需要手动加cudaDeviceSynchronize()或者用cudaStreamSynchronize()。我见过有人因为这个原因debug了三天。
动态并行适合什么场景？递归算法、树遍历、自适应网格细化。但说实话，能用静态并行解决的问题，尽量别用动态并行。为什么呢？
- 动态并行有额外的启动开销（约10μs）
- 调试困难，你没法在子核函数里打断点
- 性能不稳定，取决于GPU的调度策略
4.5 启动参数优化实战
优化启动参数，我总结了三步走：
- 确定block大小：从128开始，用profiler跑一遍，然后试256、512，看哪个occupancy最高
- 调整grid大小：保证每个SM至少有4~8个active block，这样能隐藏访存延迟
- 微调shared memory：如果用了动态共享内存，尽量让sharedMem是32的倍数，避免bank conflict
避坑指南：我曾经在优化一个图像处理核函数时，发现occupancy高达100%，但性能就是上不去。后来发现是block大小选得太小，导致每个block的计算量太少，SM的指令发射单元一直在空转。记住：occupancy高不等于性能好，还要看计算密度。
最后分享一个实用技巧：用cudaOccupancyMaxPotentialBlockSize这个API来自动计算最优block大小。虽然不能保证100%最优，但至少是个不错的起点。
int minGridSize, blockSize;
cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize,
kernel, 0, 0);
int gridSize = (N + blockSize - 1) / blockSize;嗯，启动配置这块就聊这么多。记住：没有银弹，每个核函数都有自己的最佳配置。多跑profiler，多试几次，慢慢就有感觉了。