写GPGPU的文章最怕一开始就掉进“堆术语”的坑。Microarchitecture、Vector Register、Scoreboard这些词砸出来新手直接劝退老手也觉得你只是在复读手册。这篇东西我想换个写法从“为什么需要它”讲起把GPGPU的编程模型和它背后的架构逻辑串成一条线。你不是来背名词的你是来搞懂“那玩意儿到底怎么把上万个小核捏在一起干活的”。这个系列我会拆成几篇第一篇先把骨架搭起来也就是SPMD编程模型、线程层次结构、以及它和SIMD之间的那点纠缠不清的关系。1. 为什么CPU搞不定的活要交给GPGPU先想一个问题一个4.0GHz的现代CPU核心理论上每秒能执行大概80亿次操作这算力听着挺吓人对吧但你拿它去做一张1080p图像的模糊处理哪怕是一个简单的3x3卷积核每个像素要读9个像素再算9次乘加一共1920x1080x9约等于1866万次操作单核CPU跑起来也就是个位数毫秒的级别也不算慢。可一旦你换成分辨率4K、60fps的视频实时处理每帧830万像素每秒要处理将近5亿像素每个像素再来几十次浮点运算CPU单核那80亿次操作看起来就不够分了——而且你还得同时跑操作系统、跑游戏逻辑、跑网络协议栈。这里头就是GPGPU存在的根本逻辑CPU追求的是“低延迟下的复杂逻辑处理”GPU追求的是“高吞吐下的海量并行计算”。它们是两种完全不同的设计哲学。CPU把大量的芯片面积花在分支预测器、乱序执行引擎、大容量缓存上就是为了让单条指令的执行延迟尽量低而GPGPU把这些面积几乎全砍了把省下来的晶体管全部做成ALU算术逻辑单元用“数量”来碾压“速度”。用个不太精确但好理解的类比CPU是几个博士生数学题做得飞快但一次只能做一道GPU是一万个初中生一个人做一道题速度不算快但一万人同时开算总吞吐量不是一个量级。你要做的是那种能拆成一万份的题这时候GPU的优势就体现出来了。所以GPGPU真正解决的问题是在可控的功耗和成本范围内为数据并行Data Parallelism类型的计算提供远超CPU的算力。图像处理、物理仿真、深度学习训练、科学计算里的矩阵运算全是这类——数据量大逻辑相对统一每条数据执行的操作几乎一样。你要是让GPU去跑那种分支极度复杂、数据依赖极强的串行逻辑它能被CPU按在地上摩擦。2. GPGPU的核心抽象SIMT与SPMD搞懂GPGPU必须跨过的一道坎就是分辨三个长得特别像的词SIMD单指令多数据、SPMD单程序多数据、SIMT单指令多线程。CPU的SIMD比如AVX-512是处理器内有一组宽向量寄存器一条指令能同时对8个float做加法。这是硬件层面一次性把多条数据塞进一条指令里算。SPMD是软件层面的写法你写一份代码按逻辑它会被多个处理单元同时执行每个处理单元处理自己的那部分数据。你写CUDA或者OpenCL的kernel的时候其实就是SPMD思路——一份kernel函数成千上万个线程同时去跑这份代码但每个线程通过自己的全局ID去取不同的数据来处理。GPGPU真正有点暧昧的地方在SIMT这是NVIDIA提出的一个硬件执行模型它想让你感觉自己在写多线程的SPMD代码但底层硬件是用类似SIMD的方式成组执行线程的。这就是GPGPU和CPU在设计理念上最重要的分岔路之一CPU是硬件帮你调度线程你尽管创建线程就行操作系统和CPU核心帮你切换GPU是硬件把一堆线程捆绑在一起以固定的步调同步推进。这个“捆绑”的单位NVIDIA叫Warp一般是32个线程AMD叫Wavefront一般是64个线程。所以你在写GPGPU程序时脑子里要时刻绷着这根弦**我写的每一行“看似是单线程逻辑”的代码实际上是被硬件打包成Warp后以锁步lock-step方式执行的。**这直接引出了后面无穷无尽的性能陷阱最典型的就是“分支发散”——同一Warp里的线程因为数据不同走到了一个if的两个分支里硬件没办法同时执行两个分支只能把两个分支都执行一遍没走那个分支的线程就白等着性能直接腰斩甚至更惨。3. 编程模型与架构原理解析3.1 线程层次Grid、Block、Warp写CUDA代码你第一眼看到的就是那套奇怪的线程层次结构。它不只是软件概念而是精准对应了硬件组织结构。从最上层往下拆Grid你启动一个kernel时创建的整个线程集合Block又称CTACooperative Thread Array协作线程数组Grid里分成若干组ThreadBlock里单个执行流硬件上是这么对应的一个Block会被调度到一个SMStreaming Multiprocessor流式多处理器上执行一个SM会把Block里的线程再次切成一个个Warp来实际执行。为什么中间要多一层Block这层抽象本质上是为了让线程之间可以协作。同一个Block内的线程可以通过共享内存Shared Memory交换数据还可以通过屏障同步Barrier来对齐执行进度。不同Block之间是不允许直接同步的因为硬件不保证它们会同时运行甚至可能一个在跑一个还在排队。写代码时怎么定这些参数经验法则Block大小建议取32的整数倍因为Warp大小是32。取64、128、256都是常见的配置。Block太小比如32可能让一个SM上的线程数太少不足以藏延迟Block太大比如1024又受限于每个SM的最大线程数通常是1024或2048容易导致一个SM只能放下一个Block调度灵活性下降。3.2 存储层次谁快谁慢你得心里有数GPGPU最劝退新手的地方就是它的存储层次比CPU复杂得多而且每一层的命中和不命中代价差距是数量级的。从快到慢大致排寄存器文件Register File每线程私有最快。SM内部就是一大块寄存器堆一个SM可能有65536个32位寄存器分给每个线程的寄存器数是可以配置的。寄存器溢出Spill是性能杀手因为溢出去的会跑到本地内存而本地内存实际上在显存里。共享内存Shared Memory同一Block内线程共享在SM芯片上延迟大概是全局内存的十分之一上下带宽极高。这是GPGPU优化里最常抠的地方矩阵分块乘法的核心就是把数据搬进共享内存反复重用。只读/纹理缓存Read-Only/Texture CacheGPU有专门的纹理硬件单元除了图像采样它也可以用来当只读缓存对某些访问模式有奇效。L1/L2缓存L1在每个SM内部L2在整个GPU内共享。注意CPU上缓存一致性是个大课题GPU上这块相对简化因为GPU假设你用的是“直通”模型——直接读全局内存数据是一致的一旦你用了共享内存或者做了Atomic操作那就得额外小心。全局内存Global Memory就是显存几GB到几十GB。带宽高但延迟可能几百个周期。GPGPU性能优化里最核心的一条定律尽量减少全局内存的访问次数尽量让访问是合并的Coalesced。主机内存Host Memory通过PCIe总线连接带宽和延迟跟显存完全不在一个量级。这决定了你在CPU和GPU之间拷贝数据时要格外抠门能不拷就不拷能一次性拷完就不要分批拷。3.3 Warp执行与延迟隐藏前面提到Warp是32个线程成组执行那这一组线程到底是怎么跑的更准确地说一个Warp里的32个线程如果走的是同一条路径那它们执行一条指令的时间和一个线程执行一条指令差不多。SM里实际的计算单元CUDA Core数量是有限的比如常见的一个SM有128个CUDA Core。一块卡可能有好几千个CUDA Core但当Warp数量远超物理核心数量时SM用时间片轮转的方式让Warp们轮流上计算单元执行。这背后的核心机制叫延迟隐藏Latency Hiding。比如一条全局内存load指令发出后可能要400个周期才能拿到数据。CPU的做法是等或者做分支预测、乱序执行来填坑GPU的做法简单粗暴——切换到另一个Warp去执行反正SM上有几十个Warp等着呢。GPU就是要让你同时有足够多的Warp在跑让内存延迟被计算掩盖住。这个道理深刻影响了编程模型在GPGPU里写出大量可以并行执行的轻量线程不是“给硬件找麻烦”而是“给硬件喂饭吃”。你的线程越多硬件越容易找到Warp来切换延迟隐藏效果越好。反过来如果你只启动了和CUDA Core数量相等的线程那硬件没得切每个线程都得裸着等内存延迟性能可能惨不忍睹。4. 实操一个最简单的GPGPU程序从头跑通理论说了那么多不写代码等于白说。下面用一个最简的CUDA例子带你过一遍GPGPU编程的基本流程。虽然语言是CUDA但核心思路在OpenCL、HIP、SYCL里完全通用。4.1 环境准备我用的环境是Ubuntu 22.04一张NVIDIA RTX 3060Ampere架构8组GPC28个SMCUDA Toolkit 12.x。装好之后命令行输入nvcc --version能正常输出版本号就行。如果你手头没有NVIDIA显卡也没关系直接用Google Colab的免费GPU或者用AMD的ROCm平台跑HIP版本代码逻辑完全一致。4.2 第一个Kernel向量相加我们要做的是A、B两个数组对应元素相加结果存入C。这是GPGPU界的“Hello World”但别小看它它把GPGPU最关键的几个要素全涉及了显存分配、数据拷贝、Kernel启动、线程索引计算。#include cstdio // GPU端执行的函数前面加 __global__ 修饰符 __global__ void vectorAdd(const float* a, const float* b, float* c, int n) { // 计算当前线程对应的全局索引 int idx blockIdx.x * blockDim.x threadIdx.x; if (idx n) { c[idx] a[idx] b[idx]; } } int main() { int n 1 20; // 1048576个元素 size_t bytes n * sizeof(float); // 1. 在GPU上分配显存 float *d_a, *d_b, *d_c; cudaMalloc(d_a, bytes); cudaMalloc(d_b, bytes); cudaMalloc(d_c, bytes); // 2. 准备主机端数据 float *h_a (float*)malloc(bytes); float *h_b (float*)malloc(bytes); float *h_c (float*)malloc(bytes); for (int i 0; i n; i) { h_a[i] 1.0f; h_b[i] 2.0f; } // 3. 把数据从主机内存拷贝到显存 cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice); cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice); // 4. 启动Kernel配置Grid和Block大小 int blockSize 256; int gridSize (n blockSize - 1) / blockSize; vectorAddgridSize, blockSize(d_a, d_b, d_c, n); // 5. 把结果拷贝回主机 cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost); // 6. 验证结果 for (int i 0; i n; i) { if (h_c[i] ! 3.0f) { printf(Error at index %d: %f\n, i, h_c[i]); break; } } printf(Done! Sample h_c[0] %f\n, h_c[0]); // 7. 善后释放显存和内存 cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); free(h_a); free(h_b); free(h_c); return 0; }编译运行nvcc vector_add.cu -o vector_add ./vector_add如果输出Done! Sample h_c[0] 3.000000说明整个流程跑通了。4.3 这段代码背后发生了什么一步步拆解一下因为很多新手虽然把代码跑通了但根本不知道硬件上发生了什么。cudaMalloc是在显存上分配内存返回的是设备指针。注意这个指针你不能在主机端直接解引用必须通过cudaMemcpy来传数据。vectorAddgridSize, blockSize这套尖括号语法是CUDA启动Kernel的特殊写法第一个参数是Grid里Block的数量第二个参数是每个Block里Thread的数量。这里gridSize算出来是4096blockSize是256总线程数是1048576正好等于n所以不需要越界判断也能全部覆盖但为了通用性我加了if (idx n)来判断这样即使n不是blockSize的整数倍也不会越界访问。在kernel函数体里blockIdx.x * blockDim.x threadIdx.x这个索引计算公式是GPGPU编程的基础功。它的意思是先算出当前线程在自己的Block里的位置threadIdx.x再算出当前Block在Grid里是第几个blockIdx.x乘以每个Block的线程数blockDim.x就得到全局唯一索引。硬件执行时是这样的4096, 256启动后这4096个Block会被分发到GPU的28个SM上。每个SM分到的Block数量不固定看调度器的决定。每当把一个Block分配到SM上SM会把Block里的256个线程切成8个Warp256 / 32 8然后每个Warp一个接一个地执行。这里有个值得注意的细节Block里的线程是顺序执行还是并行执行其实不完全由你控制。一个SM如果有足够的计算单元可能多个Warp同时占用如果资源不够就得排队。但这些对你是透明的你只需要保证Block内部的线程数量配置合理就行。4.4 OpenCL版本对照很多人可能不是NVIDIA平台或者想写跨厂商的代码这里给一个OpenCL版本的对照让大家理解CUDA和OpenCL在概念上的映射关系。OpenCL的代码模板化程度更高写起来更啰嗦但底层的概念完全一致。// cl_kernels.cl __kernel void vectorAdd(__global const float* a, __global const float* b, __global float* c, const unsigned int n) { int idx get_global_id(0); if (idx n) { c[idx] a[idx] b[idx]; } }主机的OpenCL代码会比CUDA啰嗦很多因为你需要显式管理平台、设备、上下文、命令队列、程序对象但核心的4步是一样的分配设备内存 - 写数据到设备 - 启动Kernel - 读回结果。这里不再展开完整代码大家知道概念映射关系即可__global对应CUDA的全局内存指针get_global_id(0)对应blockIdx.x * blockDim.x threadIdx.x。5. 常见误区与性能排查实录5.1 误区一线程数越多越好很多人以为GPU不就是核多吗那我就使劲开线程开它个上千万个。方向是对的但要注意度。线程数太多会导致每个线程分到的寄存器数太少或者Block数太多导致调度开销增大反而性能下降。经验是总线程数覆盖数据规模是底线Block大小在128-512之间通常是最优区间。对于现代NVIDIA架构每个SM能同时驻留的最大线程数一般是1536或2048。你可以算一下比如RTX 3060有28个SM每个SM最大驻留1536个线程那一共能同时驻留43008个线程。你启动一百万线程意味着大部分时间在排队。这没问题GPU就是靠“排队切换”来隐藏延迟的。但如果你一个Warp里只有一个线程在干正经事那才是灾难。一个真正常见的问题某个Block的执行时间很长但Block内部一开始就把大部分线程阻塞在一个__syncthreads()上剩下的线程慢慢吞吞地执行。这就把整个Block的进度拖慢了进而拖慢整个Grid。所以__syncthreads()不是随便用的每用一次Block里所有线程都得等对齐这是硬同步。5.2 误区二网格大小随便取整就行很多新手写int gridSize n / blockSize;然后发现当n不能被blockSize整除时最后一部分数据没被处理。其实这个问题很简单用向上取整就行int gridSize (n blockSize - 1) / blockSize;然后在kernel里加越界判断。但就是这个小问题在一些生产代码里能看到好多次。尤其在做科学计算时数据量N不是2的幂次的情况太常见了。另外提一下CUDA 6.0以后引入了CUDA Graphs技术可以把多个kernel的启动图记录下来一次性重放。在kernel启动次数非常多比如几万次的时候这个能把启动开销从微秒级压到纳秒级。这是优化大循环里kernel重复启动的好办法但现在很多教程里都没提到值得自己去看看官方文档。5.3 误区三显存拷贝不花钱我见过有人每次迭代都在CPU和GPU之间拷贝整个数据集理由是“方便调试”。这种代码的性能基本是灾难级的。PCIe 4.0 x16的理论带宽是32GB/s这和GPU显存动辄几百GB/s的带宽差了快一个数量级。正确的做法是能留在GPU上就留在GPU上只在必要时才和主机通信。在现代CUDA里还有统一虚拟地址空间Unified Virtual Address和零拷贝内存Zero-Copy Memory可以用但它们的性能特性跟显存不同你得想清楚再用。如果非要频繁小批量拷贝也有一个经验合并小拷贝为批量异步拷贝。cudaMemcpyAsync配合cudaStream可以实现拷贝和计算的重叠这个在流水线优化里是杀手锏比单纯减少拷贝次数更不吃代码结构的亏。5.4 用Profiler定位性能瓶颈而不是猜性能不好第一反应不应该是“我猜是这里慢”而是上工具。NVIDIA的ncuNsight Compute是看kernel内部性能的利器ncu --set full能给出详细的执行资源利用率、内存吞吐、Warp状态、分支效率等数据。举一个我自己的例子。之前做一个粒子碰撞模拟GPU占用率看着有90%但速度就是上不去。用ncu一测发现Stall Long Scoreboard占了将近一半的Warp停顿时间说明大量Warp在等内存数据。再一看内存访问模式发现一个数组按列访问了而数组是行主序存储在显存里的每次读取整行却只用其中一个元素导致内存带宽浪费了4倍。改成结构体数组SoA布局之后速度直接提升了1.6倍。很多时候性能问题不是靠“优化代码逻辑”解决的而是靠调整数据布局解决的。这是GPGPU优化和CPU优化一个很大的不同点CPU上缓存不友好可能只慢10%GPU上全局内存访问模式不对可能慢好几倍。因为GPU的算力密度实在太夸张喂不饱它它就闲着。6. 后续内容预告与个人经验谈这个系列如果只是讲“怎么把CUDA跑起来”那和看官方手册区别不大。我打算接下来用一篇专门讲内存访问模式与优化实战从合并访问Coalesced Access、共享内存的Bank Conflict到矩阵乘法的分块优化一步步把性能从最开始的实现往上拉几个量级。那种“同一份算法只是改改数据访问方式性能翻好几倍”的案例比任何理论都有说服力。另外提一句最近在折腾Kafka的时候感触很深。Kafka那种分区、副本、顺序读写的设计和GPGPU里“高吞吐优先于低延迟”的思路本质上是一类哲学如果你不能减少总工作量那就让每一步都流水线化别有气泡。在GPGPU里你要让Warp之间没有气泡在Kafka里你要让Segment日志一直在顺序写道理是相通的。底层系统设计做到最后都是在解决“怎么让硬件一直处于忙碌状态”这一件事。回到GPGPU学习这件事上。我个人经验是最有效的学习方式不是刷文档而是把一个小算法反复优化。向量加法太小了体现不出差距建议直接上手矩阵乘法从最naive的版本开始然后一步步加共享内存、分块、向量化加载、避免Bank Conflict每一步都能看到明确的性能数字变化。这种反馈循环比读十篇教程都长本事。下一篇我会重点讲共享内存和Bank Conflict这也是GPGPU优化里最考验功力的地方。学完那篇你再回来看自己几个月前写的代码大概率会有想重写的冲动——这是好事说明你真正入了门。 GPGPU编程模型与架构原理解析：从SPMD到Warp的并行计算
写GPGPU的文章最怕一开始就掉进“堆术语”的坑。Microarchitecture、Vector Register、Scoreboard这些词砸出来新手直接劝退老手也觉得你只是在复读手册。这篇东西我想换个写法从“为什么需要它”讲起把GPGPU的编程模型和它背后的架构逻辑串成一条线。你不是来背名词的你是来搞懂“那玩意儿到底怎么把上万个小核捏在一起干活的”。这个系列我会拆成几篇第一篇先把骨架搭起来也就是SPMD编程模型、线程层次结构、以及它和SIMD之间的那点纠缠不清的关系。1. 为什么CPU搞不定的活要交给GPGPU先想一个问题一个4.0GHz的现代CPU核心理论上每秒能执行大概80亿次操作这算力听着挺吓人对吧但你拿它去做一张1080p图像的模糊处理哪怕是一个简单的3x3卷积核每个像素要读9个像素再算9次乘加一共1920x1080x9约等于1866万次操作单核CPU跑起来也就是个位数毫秒的级别也不算慢。可一旦你换成分辨率4K、60fps的视频实时处理每帧830万像素每秒要处理将近5亿像素每个像素再来几十次浮点运算CPU单核那80亿次操作看起来就不够分了——而且你还得同时跑操作系统、跑游戏逻辑、跑网络协议栈。这里头就是GPGPU存在的根本逻辑CPU追求的是“低延迟下的复杂逻辑处理”GPU追求的是“高吞吐下的海量并行计算”。它们是两种完全不同的设计哲学。CPU把大量的芯片面积花在分支预测器、乱序执行引擎、大容量缓存上就是为了让单条指令的执行延迟尽量低而GPGPU把这些面积几乎全砍了把省下来的晶体管全部做成ALU算术逻辑单元用“数量”来碾压“速度”。用个不太精确但好理解的类比CPU是几个博士生数学题做得飞快但一次只能做一道GPU是一万个初中生一个人做一道题速度不算快但一万人同时开算总吞吐量不是一个量级。你要做的是那种能拆成一万份的题这时候GPU的优势就体现出来了。所以GPGPU真正解决的问题是在可控的功耗和成本范围内为数据并行Data Parallelism类型的计算提供远超CPU的算力。图像处理、物理仿真、深度学习训练、科学计算里的矩阵运算全是这类——数据量大逻辑相对统一每条数据执行的操作几乎一样。你要是让GPU去跑那种分支极度复杂、数据依赖极强的串行逻辑它能被CPU按在地上摩擦。2. GPGPU的核心抽象SIMT与SPMD搞懂GPGPU必须跨过的一道坎就是分辨三个长得特别像的词SIMD单指令多数据、SPMD单程序多数据、SIMT单指令多线程。CPU的SIMD比如AVX-512是处理器内有一组宽向量寄存器一条指令能同时对8个float做加法。这是硬件层面一次性把多条数据塞进一条指令里算。SPMD是软件层面的写法你写一份代码按逻辑它会被多个处理单元同时执行每个处理单元处理自己的那部分数据。你写CUDA或者OpenCL的kernel的时候其实就是SPMD思路——一份kernel函数成千上万个线程同时去跑这份代码但每个线程通过自己的全局ID去取不同的数据来处理。GPGPU真正有点暧昧的地方在SIMT这是NVIDIA提出的一个硬件执行模型它想让你感觉自己在写多线程的SPMD代码但底层硬件是用类似SIMD的方式成组执行线程的。这就是GPGPU和CPU在设计理念上最重要的分岔路之一CPU是硬件帮你调度线程你尽管创建线程就行操作系统和CPU核心帮你切换GPU是硬件把一堆线程捆绑在一起以固定的步调同步推进。这个“捆绑”的单位NVIDIA叫Warp一般是32个线程AMD叫Wavefront一般是64个线程。所以你在写GPGPU程序时脑子里要时刻绷着这根弦**我写的每一行“看似是单线程逻辑”的代码实际上是被硬件打包成Warp后以锁步lock-step方式执行的。**这直接引出了后面无穷无尽的性能陷阱最典型的就是“分支发散”——同一Warp里的线程因为数据不同走到了一个if的两个分支里硬件没办法同时执行两个分支只能把两个分支都执行一遍没走那个分支的线程就白等着性能直接腰斩甚至更惨。3. 编程模型与架构原理解析3.1 线程层次Grid、Block、Warp写CUDA代码你第一眼看到的就是那套奇怪的线程层次结构。它不只是软件概念而是精准对应了硬件组织结构。从最上层往下拆Grid你启动一个kernel时创建的整个线程集合Block又称CTACooperative Thread Array协作线程数组Grid里分成若干组ThreadBlock里单个执行流硬件上是这么对应的一个Block会被调度到一个SMStreaming Multiprocessor流式多处理器上执行一个SM会把Block里的线程再次切成一个个Warp来实际执行。为什么中间要多一层Block这层抽象本质上是为了让线程之间可以协作。同一个Block内的线程可以通过共享内存Shared Memory交换数据还可以通过屏障同步Barrier来对齐执行进度。不同Block之间是不允许直接同步的因为硬件不保证它们会同时运行甚至可能一个在跑一个还在排队。写代码时怎么定这些参数经验法则Block大小建议取32的整数倍因为Warp大小是32。取64、128、256都是常见的配置。Block太小比如32可能让一个SM上的线程数太少不足以藏延迟Block太大比如1024又受限于每个SM的最大线程数通常是1024或2048容易导致一个SM只能放下一个Block调度灵活性下降。3.2 存储层次谁快谁慢你得心里有数GPGPU最劝退新手的地方就是它的存储层次比CPU复杂得多而且每一层的命中和不命中代价差距是数量级的。从快到慢大致排寄存器文件Register File每线程私有最快。SM内部就是一大块寄存器堆一个SM可能有65536个32位寄存器分给每个线程的寄存器数是可以配置的。寄存器溢出Spill是性能杀手因为溢出去的会跑到本地内存而本地内存实际上在显存里。共享内存Shared Memory同一Block内线程共享在SM芯片上延迟大概是全局内存的十分之一上下带宽极高。这是GPGPU优化里最常抠的地方矩阵分块乘法的核心就是把数据搬进共享内存反复重用。只读/纹理缓存Read-Only/Texture CacheGPU有专门的纹理硬件单元除了图像采样它也可以用来当只读缓存对某些访问模式有奇效。L1/L2缓存L1在每个SM内部L2在整个GPU内共享。注意CPU上缓存一致性是个大课题GPU上这块相对简化因为GPU假设你用的是“直通”模型——直接读全局内存数据是一致的一旦你用了共享内存或者做了Atomic操作那就得额外小心。全局内存Global Memory就是显存几GB到几十GB。带宽高但延迟可能几百个周期。GPGPU性能优化里最核心的一条定律尽量减少全局内存的访问次数尽量让访问是合并的Coalesced。主机内存Host Memory通过PCIe总线连接带宽和延迟跟显存完全不在一个量级。这决定了你在CPU和GPU之间拷贝数据时要格外抠门能不拷就不拷能一次性拷完就不要分批拷。3.3 Warp执行与延迟隐藏前面提到Warp是32个线程成组执行那这一组线程到底是怎么跑的更准确地说一个Warp里的32个线程如果走的是同一条路径那它们执行一条指令的时间和一个线程执行一条指令差不多。SM里实际的计算单元CUDA Core数量是有限的比如常见的一个SM有128个CUDA Core。一块卡可能有好几千个CUDA Core但当Warp数量远超物理核心数量时SM用时间片轮转的方式让Warp们轮流上计算单元执行。这背后的核心机制叫延迟隐藏Latency Hiding。比如一条全局内存load指令发出后可能要400个周期才能拿到数据。CPU的做法是等或者做分支预测、乱序执行来填坑GPU的做法简单粗暴——切换到另一个Warp去执行反正SM上有几十个Warp等着呢。GPU就是要让你同时有足够多的Warp在跑让内存延迟被计算掩盖住。这个道理深刻影响了编程模型在GPGPU里写出大量可以并行执行的轻量线程不是“给硬件找麻烦”而是“给硬件喂饭吃”。你的线程越多硬件越容易找到Warp来切换延迟隐藏效果越好。反过来如果你只启动了和CUDA Core数量相等的线程那硬件没得切每个线程都得裸着等内存延迟性能可能惨不忍睹。4. 实操一个最简单的GPGPU程序从头跑通理论说了那么多不写代码等于白说。下面用一个最简的CUDA例子带你过一遍GPGPU编程的基本流程。虽然语言是CUDA但核心思路在OpenCL、HIP、SYCL里完全通用。4.1 环境准备我用的环境是Ubuntu 22.04一张NVIDIA RTX 3060Ampere架构8组GPC28个SMCUDA Toolkit 12.x。装好之后命令行输入nvcc --version能正常输出版本号就行。如果你手头没有NVIDIA显卡也没关系直接用Google Colab的免费GPU或者用AMD的ROCm平台跑HIP版本代码逻辑完全一致。4.2 第一个Kernel向量相加我们要做的是A、B两个数组对应元素相加结果存入C。这是GPGPU界的“Hello World”但别小看它它把GPGPU最关键的几个要素全涉及了显存分配、数据拷贝、Kernel启动、线程索引计算。#include cstdio // GPU端执行的函数前面加 __global__ 修饰符 __global__ void vectorAdd(const float* a, const float* b, float* c, int n) { // 计算当前线程对应的全局索引 int idx blockIdx.x * blockDim.x threadIdx.x; if (idx n) { c[idx] a[idx] b[idx]; } } int main() { int n 1 20; // 1048576个元素 size_t bytes n * sizeof(float); // 1. 在GPU上分配显存 float *d_a, *d_b, *d_c; cudaMalloc(d_a, bytes); cudaMalloc(d_b, bytes); cudaMalloc(d_c, bytes); // 2. 准备主机端数据 float *h_a (float*)malloc(bytes); float *h_b (float*)malloc(bytes); float *h_c (float*)malloc(bytes); for (int i 0; i n; i) { h_a[i] 1.0f; h_b[i] 2.0f; } // 3. 把数据从主机内存拷贝到显存 cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice); cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice); // 4. 启动Kernel配置Grid和Block大小 int blockSize 256; int gridSize (n blockSize - 1) / blockSize; vectorAddgridSize, blockSize(d_a, d_b, d_c, n); // 5. 把结果拷贝回主机 cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost); // 6. 验证结果 for (int i 0; i n; i) { if (h_c[i] ! 3.0f) { printf(Error at index %d: %f\n, i, h_c[i]); break; } } printf(Done! Sample h_c[0] %f\n, h_c[0]); // 7. 善后释放显存和内存 cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); free(h_a); free(h_b); free(h_c); return 0; }编译运行nvcc vector_add.cu -o vector_add ./vector_add如果输出Done! Sample h_c[0] 3.000000说明整个流程跑通了。4.3 这段代码背后发生了什么一步步拆解一下因为很多新手虽然把代码跑通了但根本不知道硬件上发生了什么。cudaMalloc是在显存上分配内存返回的是设备指针。注意这个指针你不能在主机端直接解引用必须通过cudaMemcpy来传数据。vectorAddgridSize, blockSize这套尖括号语法是CUDA启动Kernel的特殊写法第一个参数是Grid里Block的数量第二个参数是每个Block里Thread的数量。这里gridSize算出来是4096blockSize是256总线程数是1048576正好等于n所以不需要越界判断也能全部覆盖但为了通用性我加了if (idx n)来判断这样即使n不是blockSize的整数倍也不会越界访问。在kernel函数体里blockIdx.x * blockDim.x threadIdx.x这个索引计算公式是GPGPU编程的基础功。它的意思是先算出当前线程在自己的Block里的位置threadIdx.x再算出当前Block在Grid里是第几个blockIdx.x乘以每个Block的线程数blockDim.x就得到全局唯一索引。硬件执行时是这样的4096, 256启动后这4096个Block会被分发到GPU的28个SM上。每个SM分到的Block数量不固定看调度器的决定。每当把一个Block分配到SM上SM会把Block里的256个线程切成8个Warp256 / 32 8然后每个Warp一个接一个地执行。这里有个值得注意的细节Block里的线程是顺序执行还是并行执行其实不完全由你控制。一个SM如果有足够的计算单元可能多个Warp同时占用如果资源不够就得排队。但这些对你是透明的你只需要保证Block内部的线程数量配置合理就行。4.4 OpenCL版本对照很多人可能不是NVIDIA平台或者想写跨厂商的代码这里给一个OpenCL版本的对照让大家理解CUDA和OpenCL在概念上的映射关系。OpenCL的代码模板化程度更高写起来更啰嗦但底层的概念完全一致。// cl_kernels.cl __kernel void vectorAdd(__global const float* a, __global const float* b, __global float* c, const unsigned int n) { int idx get_global_id(0); if (idx n) { c[idx] a[idx] b[idx]; } }主机的OpenCL代码会比CUDA啰嗦很多因为你需要显式管理平台、设备、上下文、命令队列、程序对象但核心的4步是一样的分配设备内存 - 写数据到设备 - 启动Kernel - 读回结果。这里不再展开完整代码大家知道概念映射关系即可__global对应CUDA的全局内存指针get_global_id(0)对应blockIdx.x * blockDim.x threadIdx.x。5. 常见误区与性能排查实录5.1 误区一线程数越多越好很多人以为GPU不就是核多吗那我就使劲开线程开它个上千万个。方向是对的但要注意度。线程数太多会导致每个线程分到的寄存器数太少或者Block数太多导致调度开销增大反而性能下降。经验是总线程数覆盖数据规模是底线Block大小在128-512之间通常是最优区间。对于现代NVIDIA架构每个SM能同时驻留的最大线程数一般是1536或2048。你可以算一下比如RTX 3060有28个SM每个SM最大驻留1536个线程那一共能同时驻留43008个线程。你启动一百万线程意味着大部分时间在排队。这没问题GPU就是靠“排队切换”来隐藏延迟的。但如果你一个Warp里只有一个线程在干正经事那才是灾难。一个真正常见的问题某个Block的执行时间很长但Block内部一开始就把大部分线程阻塞在一个__syncthreads()上剩下的线程慢慢吞吞地执行。这就把整个Block的进度拖慢了进而拖慢整个Grid。所以__syncthreads()不是随便用的每用一次Block里所有线程都得等对齐这是硬同步。5.2 误区二网格大小随便取整就行很多新手写int gridSize n / blockSize;然后发现当n不能被blockSize整除时最后一部分数据没被处理。其实这个问题很简单用向上取整就行int gridSize (n blockSize - 1) / blockSize;然后在kernel里加越界判断。但就是这个小问题在一些生产代码里能看到好多次。尤其在做科学计算时数据量N不是2的幂次的情况太常见了。另外提一下CUDA 6.0以后引入了CUDA Graphs技术可以把多个kernel的启动图记录下来一次性重放。在kernel启动次数非常多比如几万次的时候这个能把启动开销从微秒级压到纳秒级。这是优化大循环里kernel重复启动的好办法但现在很多教程里都没提到值得自己去看看官方文档。5.3 误区三显存拷贝不花钱我见过有人每次迭代都在CPU和GPU之间拷贝整个数据集理由是“方便调试”。这种代码的性能基本是灾难级的。PCIe 4.0 x16的理论带宽是32GB/s这和GPU显存动辄几百GB/s的带宽差了快一个数量级。正确的做法是能留在GPU上就留在GPU上只在必要时才和主机通信。在现代CUDA里还有统一虚拟地址空间Unified Virtual Address和零拷贝内存Zero-Copy Memory可以用但它们的性能特性跟显存不同你得想清楚再用。如果非要频繁小批量拷贝也有一个经验合并小拷贝为批量异步拷贝。cudaMemcpyAsync配合cudaStream可以实现拷贝和计算的重叠这个在流水线优化里是杀手锏比单纯减少拷贝次数更不吃代码结构的亏。5.4 用Profiler定位性能瓶颈而不是猜性能不好第一反应不应该是“我猜是这里慢”而是上工具。NVIDIA的ncuNsight Compute是看kernel内部性能的利器ncu --set full能给出详细的执行资源利用率、内存吞吐、Warp状态、分支效率等数据。举一个我自己的例子。之前做一个粒子碰撞模拟GPU占用率看着有90%但速度就是上不去。用ncu一测发现Stall Long Scoreboard占了将近一半的Warp停顿时间说明大量Warp在等内存数据。再一看内存访问模式发现一个数组按列访问了而数组是行主序存储在显存里的每次读取整行却只用其中一个元素导致内存带宽浪费了4倍。改成结构体数组SoA布局之后速度直接提升了1.6倍。很多时候性能问题不是靠“优化代码逻辑”解决的而是靠调整数据布局解决的。这是GPGPU优化和CPU优化一个很大的不同点CPU上缓存不友好可能只慢10%GPU上全局内存访问模式不对可能慢好几倍。因为GPU的算力密度实在太夸张喂不饱它它就闲着。6. 后续内容预告与个人经验谈这个系列如果只是讲“怎么把CUDA跑起来”那和看官方手册区别不大。我打算接下来用一篇专门讲内存访问模式与优化实战从合并访问Coalesced Access、共享内存的Bank Conflict到矩阵乘法的分块优化一步步把性能从最开始的实现往上拉几个量级。那种“同一份算法只是改改数据访问方式性能翻好几倍”的案例比任何理论都有说服力。另外提一句最近在折腾Kafka的时候感触很深。Kafka那种分区、副本、顺序读写的设计和GPGPU里“高吞吐优先于低延迟”的思路本质上是一类哲学如果你不能减少总工作量那就让每一步都流水线化别有气泡。在GPGPU里你要让Warp之间没有气泡在Kafka里你要让Segment日志一直在顺序写道理是相通的。底层系统设计做到最后都是在解决“怎么让硬件一直处于忙碌状态”这一件事。回到GPGPU学习这件事上。我个人经验是最有效的学习方式不是刷文档而是把一个小算法反复优化。向量加法太小了体现不出差距建议直接上手矩阵乘法从最naive的版本开始然后一步步加共享内存、分块、向量化加载、避免Bank Conflict每一步都能看到明确的性能数字变化。这种反馈循环比读十篇教程都长本事。下一篇我会重点讲共享内存和Bank Conflict这也是GPGPU优化里最考验功力的地方。学完那篇你再回来看自己几个月前写的代码大概率会有想重写的冲动——这是好事说明你真正入了门。