CUDA并行计算优化技巧
CUDA并行计算优化技巧
- 线程 / 网格配置：别让 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>>>(…);
- 全局内存访问：一定要“顺着读、顺着写”
这是 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