【CUDA编程入门②】CUDA性能优化基础内存访问、Warp Shuffle与Occupancy调优本文是「CUDA编程与HPC高性能计算」系列的第二篇。会写CUDA kernel只是第一步写得快才是真本事。CUDA性能优化说复杂也复杂说简单也简单——核心就是两件事内存和计算。本文把最基础、最常用的优化技术一网打尽掌握这些你的kernel性能就能上一个大台阶。一、写在前面优化之前先判断瓶颈在讲优化技术之前先强调一个最重要的原则先判断瓶颈再做优化。不知道瓶颈在哪就瞎优化90%都是无用功。CUDA kernel的性能瓶颈通常分两类瓶颈类型含义优化方向Memory-Bound访存瓶颈显存带宽不够喂不饱计算单元减少访存、优化访存模式、提高数据复用Compute-Bound计算瓶颈计算单元不够用带宽还有富余提高算力利用率、用Tensor Core、减少计算量怎么判断用Nsight Computencu看显存带宽利用率接近峰值 → memory-bound计算单元SM/Tensor Core利用率接近峰值 → compute-bound或者用算术强度Arithmetic Intensity估算算术强度 FLOPs / 访存字节数算术强度低 → memory-bound算术强度高 → compute-bound。先搞清楚卡在哪再对症下药。先上一张优化技术全景图二、内存访问优化CUDA优化的第一课大部分kernel尤其是初学者写的都是memory-bound的——所以内存访问优化是重中之重。2.1 合并访问Coalesced Access这是最基础、也是性价比最高的优化。原理一个Warp32个线程如果访问一段连续的全局内存硬件可以把32个请求合并成少数几个cache line传输——带宽利用率最高。✅ 合并访问理想 Thread 0 → addr 0 Thread 1 → addr 4 Thread 2 → addr 8 ... Thread 31 → addr 124 → 32个请求合并为1-2个cache line传输 ❌ 跨步访问糟糕 Thread 0 → addr 0 Thread 1 → addr 4096 Thread 2 → addr 8192 ... → 每个线程一个cache line带宽利用率极低怎么做到合并访问让thread i访问第i个连续的元素也就是data[blockIdx.x * blockDim.x threadIdx.x]这是最自然的写法但很多人写着写着就歪了——比如按列访问矩阵、跨步访问数组经验法则90%的情况下你的kernel应该是合并访问的。如果不是先想想能不能改写成合并的。2.2 共享内存预取Shared Memory Tiling如果数据需要多次复用那每次都从全局内存读就太亏了。这时候可以全局内存 → 共享内存 → 计算 一次高延迟访问 多次低延迟访问先把数据从全局内存加载到共享内存慢但只做一次然后线程从共享内存读取快可以读很多次。典型应用场景矩阵乘法分块GEMM Tiling卷积im2colStencil计算每个点需要周围点的数据关键前提数据复用率要高。如果每个数据只用一次那搬到共享内存反而多此一举——直接合并读全局内存更快。2.3 向量化加载使用float4、int4等向量类型一次加载多个元素float4 vecreinterpret_castfloat4*(gmem)[idx];// 一次加载4个floatvec.x, vec.y, vec.z, vec.w好处减少指令数量1条指令顶4条减少访存事务数有助于编译器做更好的优化什么情况下不能用数据不是对齐的比如数组起始地址不是16字节对齐数据量不是4的倍数边界处理麻烦访问模式不连续向量化也救不了三、Warp ShuffleWarp内通信的神器3.1 什么是Warp ShuffleWarp Shuffle是一种特殊指令允许同一个Warp内的线程直接交换寄存器中的数据——不需要经过共享内存也不需要同步。你想想之前Warp内线程要交换数据怎么办线程A把数据写到共享内存__syncthreads()同步一下线程B从共享内存读出来又要共享内存又要同步麻烦还慢。有了Warp Shuffle直接在寄存器层面交换——又快又省事。3.2 常用指令__shfl_sync(mask,var,srcLane);// 从指定lane读取__shfl_up_sync(mask,var,delta);// 从 lane-delta 读取往上__shfl_down_sync(mask,var,delta);// 从 lanedelta 读取往下__shfl_xor_sync(mask,var,laneMask);// 蝴蝶交换按位异或举个例子__shfl_down_sync(0xFFFFFFFF, val, 16)lane 0 从 lane 16 读数据lane 1 从 lane 17 读数据…lane 15 从 lane 31 读数据lane 16-31 从哪里读超出范围了读不到返回原数据或未定义3.3 优势与局限优势局限不需要共享内存只能在一个Warp内通信32个线程不需要__syncthreads()跨Warp通信还是得用共享内存速度快寄存器级只能交换32-bit或64-bit的数据典型应用Reduce归约Warp内求和Scan前缀和矩阵乘法累加Warp内数据交换后面讲Reduce算子的时候我们会看到Warp Shuffle的具体用法。四、Occupancy与资源分配4.1 什么是Occupancy复习一下Occupancy 每个SM上实际驻留的Warp数 / 最大可驻留Warp数每个SM的资源是有限的以A100为例资源典型值最大驻留Warp数64最大驻留线程数2048最大共享内存164 KB寄存器文件256 KB如果每个线程用64个寄存器每个Block 256线程一个Block用的寄存器 256 × 64 16384 个 64KB一个SM最多同时驻留 256KB / 64KB 4个Block 1024线程 32个WarpOccupancy 32 / 64 50%4.2 怎么提高Occupancy减少寄存器使用编译选项--maxrregcount限制每个线程最多用多少寄存器代码层面减少局部变量、复用变量或者用__launch_bounds__提示编译器减少共享内存使用优化Tile大小用寄存器代替共享内存如果数据量小调整Block大小128或256通常是比较好的选择太小调度开销大太大资源不够用4.3 常见误区Occupancy越高越好错Occupancy高有助于隐藏内存延迟一些Warp等内存的时候其他Warp可以干活但如果kernel已经是compute-bound提高Occupancy收益很小为了提高Occupancy而强制减少寄存器可能导致寄存器溢出反而更慢有些kernel本身计算密集低Occupancy也能跑满算力结论Occupancy是参考指标不是越高越好。要结合具体kernel分析以实际性能为准。五、同步与原子操作5.1 同步原语__syncthreads();// Block内所有线程同步Barrier__syncwarp(mask);// Warp内同步__threadfence();// 全局内存顺序保证__threadfence_block();// Block内顺序保证__syncthreads()是最常用的——同Block内所有线程都走到这里才继续往下走。典型场景所有线程把数据写到共享内存后同步一下然后再读确保大家都写完了分阶段的计算确保前一阶段都完成了再进入下一阶段注意__syncthreads()必须所有线程都能到达——如果放在if分支里而且不是所有线程都走这个分支会死锁5.2 原子操作atomicAdd(addr,val);// 原子加atomicMax(addr,val);// 原子取最大值atomicCAS(addr,compare,val);// 原子比较并交换原子操作保证读-改-写是原子的——不会被其他线程打断。但是原子操作是全局序列化的——很多线程同时做原子操作会排队性能很差。怎么减少原子操作先局部规约再全局写先在共享内存/Warp内做归约得到一个部分结果再用一次原子操作写回全局内存用Warp ShuffleWarp内先规约减少原子次数多级归约先在Block内归约再做第二级归约最后只有一个结果经验法则原子操作能不用就不用能用局部归约代替就代替。尤其是大量线程同时做atomicAdd性能会掉得很惨。六、优化决策路线图最后给大家一个优化的路线图遇到慢的kernel按这个顺序来kernel慢 │ ├─ 第一步profile定位 │ ├─ 用nsys看是不是GPU在等CPU/数据 │ └─ 用ncu看哪个部分最慢 │ ├─ 第二步看访存 │ ├─ 是合并访问吗→ 不是就改成合并的 │ ├─ 数据有复用吗→ 有就用共享内存预取 │ └─ 能向量化吗→ 能就用float4等 │ ├─ 第三步看计算 │ ├─ 能用Tensor Core吗→ 能用就用后面讲 │ ├─ 有多余的计算吗→ 能省就省 │ └─ 分支发散严重吗→ 优化分支 │ ├─ 第四步看资源 │ ├─ Occupancy太低→ 减少寄存器/共享内存 │ ├─ 寄存器溢出→ 减少局部变量 │ └─ Bank Conflict→ 调整共享内存访问模式 │ └─ 第五步看同步/原子 ├─ 同步太多→ 减少同步次数 └─ 原子操作多→ 先局部规约七、总结一下今天这篇文章我们讲了CUDA性能优化的基础技术内存访问优化合并访问最重要、共享内存预取数据复用、向量化加载Warp ShuffleWarp内高效通信无需共享内存和同步Occupancy受寄存器、共享内存、Block大小限制不是越高越好同步与原子操作慎用过度同步和原子竞争会严重影响性能优化路线图先profile定位再从访存、计算、资源、同步逐步优化下一篇预告我们来实战第一个经典算子——Reduce归约。从朴素的全局原子加到共享内存树形归约再到Warp Shuffle优化一步步把Reduce做到高性能。本系列导航第1篇CUDA编程体系全景线程模型、内存层次与Warp执行机制第2篇CUDA性能优化基础内存访问、Warp Shuffle与Occupancy调优本文第3篇经典算子实战Reduce归约——从朴素到高性能第4篇经典算子实战GEMM矩阵乘法——从Tiling到Tensor Core更多内容持续更新中……
