【CUDA编程入门①】CUDA编程体系全景线程模型、内存层次与Warp执行机制本文是「CUDA编程与HPC高性能计算」系列的第一篇。从这篇开始我们正式进入GPU内核编程的世界。CUDA难吗说难也难——它有一套全新的编程模型。说简单也简单——只要你搞懂了三件事线程怎么组织、内存怎么分层、Warp怎么执行CUDA就入门了。一、写在前面CUDA不是在GPU上写C很多初学者以为CUDA就是把C代码搬到GPU上跑——这是最大的误解。CUDA与普通C最大的区别在于两点线程组织模型GPU不是一个线程在跑而是几千上万个线程同时在跑而且这些线程是按特定层级组织的内存层次结构GPU不是一块内存而是有寄存器、共享内存、全局内存等多层存储速度和容量天差地别忽略了这两点写出来的kernel可能只有CPU级别的性能——那你用GPU干嘛呢先上一张全景图建立整体印象二、CUDA线程组织模型Grid/Block/Thread三级2.1 三级结构启动一个CUDA kernel时会生成一个由线程组成的网格Grid一个Kernel的所有线程 ├── Block (0,0) Block (1,0) ... │ ├── Thread(0,0) │ ├── Thread(1,0) │ └── ... └── Block (0,1) └── ...从大到小三级别层级含义维度协作能力Grid整个kernel的所有线程最多3维gridDim.x/y/z全局协作需要原子操作或多kernelBlock一组线程最多3维blockDim.x/y/z同Block内可通过共享内存__syncthreads()协作Thread单个执行单元线程索引threadIdx独立执行最关键的一点同一个Block内的线程可以协作共享内存、同步不同Block之间不能直接协作。2.2 线程索引怎么算写CUDA kernel的第一步就是搞清楚我是哪个线程我该处理哪个数据。一维索引最简单intidxblockIdx.x*blockDim.xthreadIdx.x;就是第几块 × 每块多少线程 块内第几号线程。二维索引图像处理等场景introwblockIdx.y*blockDim.ythreadIdx.y;intcolblockIdx.x*blockDim.xthreadIdx.x;intwidthgridDim.x*blockDim.x;intidxrow*widthcol;行和列分别算然后展平成一维索引。小白提示索引计算是kernel正确性的基础。写错了索引结果肯定不对。建议刚开始学的时候画个小例子比如4个block每个block 4个线程手动算几个索引验证一下。2.3 Block大小怎么选Block的大小每个Block有多少线程不是随便选的有几个原则必须是32的倍数——因为Warp大小是32下一节讲常见选择128、256、512、1024受限于三个资源寄存器数量每个线程用多少寄存器共享内存大小每个Block用多少共享内存最大线程数通常一个Block最多1024线程256是最常用的万金油选择——平衡了资源占用和并行度。三、CUDA内存模型五层存储逐级变慢GPU的内存是分层的从快到慢、从近到远┌─────────────────────────────────────┐ │ HostCPU内存 │ 可访问通过cudaMemcpy与Device交互 ├─────────────────────────────────────┤ │ Global Memory全局内存 │ 所有线程可访问容量大延迟高几百cycle ├─────────────────────────────────────┤ │ Constant Memory常量内存 │ 只读有缓存适合广播式读取 ├─────────────────────────────────────┤ │ Shared Memory共享内存 │ 同Block内线程共享低延迟~30 cycle │ 分为32个bank │ 需要避免Bank Conflict ├─────────────────────────────────────┤ │ Registers寄存器 │ 每个线程私有最快数量有限 └─────────────────────────────────────┘我们逐层来看3.1 全局内存Global Memory所有线程都能访问容量就是显存大小比如80GB访问延迟高几百个cycle**合并访问Coalesced Access**可以高效利用带宽——相邻线程访问相邻地址时硬件会把多个请求合并成一个 全局内存是最大但最慢的一层。优化的核心就是尽量少访问全局内存尽量让数据停留在离计算更近的地方。3.2 共享内存Shared Memory同一个Block内的所有线程都能访问容量很小每个SM几十KB到一百多KB速度很快~30 cycles比全局内存快一个数量级还多按32-bit划分为32个bank——如果同一个Warp内多个线程同时访问同一个bank的不同地址会触发Bank Conflict存储体冲突访问被串行化典型用途矩阵乘法分块Tiling数据复用多次读取的数据先加载到共享内存线程间通信同Block内交换数据3.3 寄存器Registers每个线程私有其他线程看不到速度最快1个cycle数量有限——如果用太多会发生寄存器溢出多余的数据放到Local Memory其实就是全局内存的一部分性能急剧下降怎么看寄存器用了多少编译时加上--ptxas-options-v参数编译器会打印每个kernel用了多少寄存器、多少共享内存。3.4 常量内存Constant Memory只读的有常量缓存适合广播式读取同一个Warp内所有线程读同一个地址如果线程访问不同地址会串行化反而更慢四、Warp执行模型32个线程齐步走4.1 Warp是什么Warp warp是一绺的意思这里翻译成线程束是GPU调度的最小单位。一个Warp固定包含32个线程同一个Warp内的线程执行完全相同的指令线程以**lock-step锁步**方式执行——大家步调一致一起走可以这样理解GPU的调度器调度的不是单个线程而是一个个Warp。Warp才是GPU调度的基本单位。4.2 分支发散Branch Divergence既然同一个Warp内的线程执行相同的指令那如果它们走了不同的分支呢比如if(threadIdx.x%20){// 路径A偶数线程走这里}else{// 路径B奇数线程走这里}答案是硬件会先执行路径A走路径B的线程被mask掉不干活然后再执行路径B走路径A的线程被mask掉。结果虽然正确但性能下降了——因为两个路径是串行执行的。这就叫分支发散Branch Divergence。怎么避免让同一个Warp内的线程尽量走同一路径按Warp对齐分配任务用Warp内的优化技巧比如Warp Shuffle后面讲小白误区不是所有分支都会导致发散。如果整个Warp都走同一个分支就不会发散。比如if (idx N)这种边界判断通常只有边界上的那个Warp会发散大部分Warp不受影响。五、三大关键优化概念5.1 合并访问Coalesced Access全局内存访问的理想模式理想Thread i 访问地址 base i * 4 访存请求可被合并为少量cache line传输 带宽利用率高 ✅ 糟糕Thread i 访问地址 base i * 1024 每个线程触发独立cache line 带宽利用率低 ❌简单说就是相邻线程访问相邻地址硬件可以把多个请求合并带宽利用率最高。这是CUDA性能优化的第一课——写kernel的时候一定要注意访存模式。5.2 Bank Conflict存储体冲突共享内存按4字节划分为32个bank存储体bank 0存地址0,32,64…bank 1存地址1,33,65…以此类推。同一个Warp32个线程内如果多个线程同时访问同一个bank的不同地址就会发生bank conflict——访问被串行化性能下降。怎么避免加padding填充比如共享内存数组多开一列让跨一行的访问错开bank调整访问模式让同一个Warp内的线程访问不同的bank 举个例子共享内存矩阵转置时按行写入后按列读取就会有bank conflict——因为同一列的元素在同一个bank里。加一列padding就能解决。5.3 Occupancy占用率Occupancy是什么Occupancy 每个SM上实际驻留的Warp数 / 最大可驻留Warp数可以理解为SM的饱满程度——越高说明SM越忙。影响Occupancy的因素每个线程的寄存器数量用得多驻留的Warp就少每个Block的共享内存大小用得多驻留的Block就少每个Block的线程数最大并发Block数限制Occupancy越高越好吗不一定。高Occupancy有助于隐藏内存延迟——当一些Warp在等内存的时候其他Warp可以继续计算但如果kernel已经是compute-bound计算瓶颈再提高Occupancy收益不大而且为了提高Occupancy而减少寄存器使用可能反而让计算变慢结论Occupancy是重要指标但不是越高越好要结合具体kernel分析。六、Kernel Launch怎么启动一个kernel6.1 启动语法kernelgrid_dim,block_dim,shared_mem_bytes,stream(args...);四个参数grid_dimGrid的维度有多少个Blockblock_dimBlock的维度每个Block有多少线程shared_mem_bytes动态共享内存大小可选静态共享内存用__shared__声明streamCUDA Stream可选默认06.2 配置示例一维配置最常用// N个元素每个block 256线程add_kernel(N255)/256,256(d_a,d_b,d_c,N);为什么是(N 255) / 256因为要向上取整——比如N257就需要2个block。二维配置图像处理等dim3block(16,16);// 每个block 16×16 256线程dim3grid((W15)/16,(H15)/16);img_kernelgrid,block(d_img,W,H);七、常见误区误区一线程数越多越快线程太多 → 寄存器压力大 → Occupancy下降 → 性能反而下降。不是线程越多越好要根据资源限制找到甜点。误区二共享内存一定比全局内存快共享内存快的前提是bank conflict少、数据复用率高。如果每个线程只读一次数据而且访问是合并的——直接读全局内存可能更快省了搬到共享内存这一步。误区三分支发散无所谓Warp内分支发散会导致串行执行性能可能直接腰斩。写kernel时要注意避免。八、总结一下今天这篇文章我们讲了CUDA编程体系的基础线程组织模型Grid/Block/Thread三级结构线程索引计算是kernel正确性的基础内存模型全局内存、常量内存、共享内存、寄存器——从大到小从慢到快Warp执行模型Warp是最小调度单位32个线程lock-step执行分支发散会损害性能三大优化概念合并访问全局内存、避免Bank Conflict共享内存、合理OccupancyKernel Launchgrid、block、共享内存、stream四个配置参数下一篇预告我们来聊CUDA性能优化的基础技术——内存访问优化、Warp Shuffle、Occupancy调优、同步与原子操作。掌握了这些你才能写出跑得动的kernel。本系列导航第1篇CUDA编程体系全景线程模型、内存层次与Warp执行机制本文第2篇CUDA性能优化基础内存访问、Warp Shuffle与Occupancy第3篇经典算子实战Reduce归约——从朴素到高性能第4篇经典算子实战GEMM矩阵乘法——从Tiling到Tensor Core更多内容持续更新中……