CUDA并行计算优化技巧

CUDA并行计算优化技巧

1. 线程 / 网格配置：别让 SM 闲着
   1.1 线程块大小（blockDim）
   • 一般选 32 的倍数（1 个 warp = 32 线程）
   • 常用范围：128 / 256 / 512 线程/块
   • 经验：
   o 小于 64：warp 太少，不容易隐藏延迟
   o 大于 1024：不合法（硬件上限），而且寄存器/共享内存压力大
   1.2 网格大小（gridDim）
   • 让 block 数量远大于 SM 数，这样调度器可以轮转执行
   比如 GPU 有 80 个 SM，你至少扔几百个 block 比较稳。
   • 一维数据典型写法：
   • int threads = 256;
   • int blocks = (N + threads - 1) / threads;
   • kernel<<<blocks, threads>>>(…);

2. 全局内存访问：一定要“顺着读、顺着写”
   这是 CUDA 性能的大头。
   2.1 访问要“合并”（coalesced）
   • 同一个 warp（32 线程）访问连续地址，才能合并成少量大访存事务。
   • 典型模式（正确）：
   • int idx = blockIdx.x * blockDim.x + threadIdx.x;
   • out[idx] = in[idx]; // 每个线程访问 idx，相邻线程访问 idx+1、idx+2……
   • 不好的模式（严重拉跨）：
   • // stride 很大，每个线程隔很远
   • out[idx] = in[idx * stride];
   2.2 AoS → SoA：结构体改成数组形式
   • 如果你有：
   • struct Pixel { float r, g, b, a; };
   • Pixel *img; // AoS
   多个线程每次只用某个字段（比如 r），那 warp 在显存里是“跳着读”，不合并。
   • 推荐改成 SoA：
   • struc