为什么过去CPU自动加速终结了
很久以前电脑之所以变快，是因为CPU主频不断提高，同样的代码跑得更快。这时候程序员根本不用管并行、内存或架构优化，只要等下一代芯片就好了。但从二十一世纪头十几年开始，主频提高遇到物理极限，功耗和热量成了瓶颈，CPU不再大幅提高时钟，而是开始把多个核心塞到一块硅片上。这意味着如果你的代码没有分裂成可以并行执行的部分，它就不会变快。
CPU努力让你觉得它还是能高效执行单线程代码，搞预测、乱序执行和大缓存来“隐藏”等待，但本质没有突破。而GPU选择了一条不同的路：它不再试图隐藏所有等待，而是拥有大量极轻量级线程，只要一些线程因访问数据、分支或慢操作暂停，马上有其他准备好的线程来顶上。GPU之所以能快，是因为永远有工作排队，而不是单线程执行更聪明。
GPU和CPU的设计哲学不同
CPU继续保留顺序执行的假象，它核心复杂，用大缓存和优化让单线程尽可能快；GPU则接受现实，认为大量轻线程和可预测的并行结构才是未来。CUDA出现之前，用GPU算通用计算还得把问题硬套成图形渲染任务。而CUDA承认GPU是通用并行机器，并提供明确控制并行、线程组织、数据在哪里的模型。
CUDA并不是把整个程序放到GPU上运行，而是指导你把确实有并行性、值得并行的部分提炼出来，通过“kernel launch”把这部分代码送到GPU执行。CPU仍然是主程序的大脑，GPU是被你调用的工具，专门干那些可以千军万马齐奔的工作。
CUDA不是把程序搬去GPU，而是精准“外包”并行任务
很多人误以为CUDA是把整个程序扔进显卡运行。事实恰恰相反：一个典型的CUDA程序主体仍在CPU上执行，只有那些能被拆成成千上万份、彼此几乎不依赖的“数据并行”任务，才被专门“外包”给GPU处理。
这个外包动作，就是启动一个“核函数”（kernel）。
核函数像一张模板，GPU会瞬间复制出海量线程，每个线程拿着不同的索引（比如矩阵中的行列号），执行完全相同的代码。
例如，传统CPU做矩阵乘法要用两层for循环遍历i和j；而在CUDA里，这些循环消失了——每个线程天然对应一个(i, j)位置，通过 blockIdx.x * blockDim.x + threadIdx.x 这行代码就能算出自己的全局编号。
这种设计迫使程序员直面一个真相：GPU不是万能加速器，它只对“天生可并行”的问题有效。如果任务内部依赖严重，强行上GPU反而更慢。
CUDA最反直觉的一点，是它强制你手动管理CPU和GPU之间的数据搬运。
你必须先在显存中分配空间，再用 cudaMemcpy 把数据从内存拷过去，跑完核函数后再拷回来。
这看似繁琐，实则是核心设计哲学：GPU的计算单元极其廉价，但跨芯片搬运数据代价高昂。一次数据传输可能耗掉成千上万次浮点运算的时间。
因此，CUDA程序快不快，关键不在于用了多少线程，而在于“每次搬数据后，能干多少活”。如果搬一次数据只做几次加减乘除，那还不如不用GPU。真正的优化方向是提升“计算强度”——即每字节数据能支撑多少次计算。这就引出了GPU编程的黄金法则：尽可能复用已搬进显存的数据，减少重复搬运。
CUDA强迫你面对内存边界问题
CPU和GPU默认不共享内存。如果GPU要处理数据，你必须：
你先在GPU上分配设备内存，
你调用函数把数据从CPU内存拷贝到GPU，
然后启动核函数并行处理，
处理完再将结果从GPU拷回CPU。
刚开始这看起来像样板代码，但它的意义正是让你意识到内存是有成本的。GPU计算其实很便宜，但数据搬运特别贵。每一次cudaMemcpy都是明确的成本。决定CUDA是否对你的问题有用，往往取决于你在这些拷贝之间做了多少真正的计算。
GPU上计算便宜，飞快；但数据搬到GPU上花费的时间不容忽视。一个优秀的CUDA程序设计就是如何让“数据搬运的成本”被“高效的并行计算”所赚回来。
显存层级决定生死，变量放哪是性能决策
GPU内部有多种存储层级，放错地方等于自废武功。
寄存器最快但每个线程独占且总量有限；
共享内存（shared memory）稍慢但同一“线程块”（block）内所有线程共享，容量小但速度远超全局显存；
全局显存（global memory）最大最慢，所有线程都能访问。
在CUDA里，声明一个变量放在哪，不是编译器自动决定的细节，而是程序员必须主动做的性能权衡。放寄存器里快但你能用的少；放共享内存里也快但要管理；放全局内存里方便但慢。这种设计让你不得不思考数据在哪儿、什么时候用，而不是随便用。
比如矩阵乘法的经典优化“分块”（tiling）：把大矩阵切成小块，先由线程块协作一次性把一块数据从全局显存搬进共享内存，同步后反复使用这块数据做多次乘加运算，最后再写回结果。这样，原本需要反复读取全局显存的操作，变成一次读取+多次高速复用，计算强度大幅提升。这种优化不是改几行代码，而是重构算法逻辑本身。
如果你的核函数每次访问全局内存只做一点点计算，那么GPU峰值算力再高也无法被充分利用。在一个简单的矩阵乘法核里，如果每次都从全局内存取数然后马上做乘加，你实际上是在等内存比在做计算。
这就引出一个核心概念：计算与内存访问的比率。如果每个内存访问可以换来大量计算，GPU就能跑满；反之GPU多数时间浪费在等内存上。
因此GPU优化看起来不像传统“微优化算术”，而更像重构算法以提高数据重用率，让每一次昂贵的内存访问能产出更多计算。
分块化和线程合作模式
减少全局内存访问最有效的办法是让线程合作重用数据。共享内存比全局内存快得多，但容量小。
所谓“分块化模式”(tiling)就是：
- 先把大数据分成合适大小的小块能装进共享内存，
- 然后一个线程块里的所有线程一起把这一块从全局内存搬到共享内存，
- 再同步，所有线程从共享内存多次读取数据做计算，
- 完成后再搬下一块。
这意味着一次全局内存访问可以被多次重用，极大提升计算与内存比。
占用率和资源竞争
共享内存和寄存器这些快存储是有限的。每个线程和每个线程块都占用一定量。如果你给每个线程太多寄存器或每个块太多共享内存，那么设备上同时能运行的线程块就变少，这反而降低并行度。
这是一种很有特色的权衡：单个线程做得更好不一定整体性能更高。GPU优化不仅看单线程如何，还要平衡全设备的线程数量与资源分配。
线程调度暗藏玄机，占不满资源等于浪费算力
GPU的计算单元叫流式多处理器（SM），每个SM能同时容纳多个“线程束”（warp）——通常32个线程为一束。
一个warp是由固定数量的线程组成的执行单元，多个warp驻留在同一个流多处理器上（SM）。GPU希望有足够的线程占满每个SM，这样才不会在某些线程等待时整体性能受损。这解释了为什么线程数量、块大小、占用率能显著影响性能。
当某个warp因读内存而卡住时，SM会零成本切换到另一个就绪的warp继续执行，实现“延迟隐藏”。这种“切换”几乎没有代价。这样即使各种等待存在，GPU整体也不会闲着。
但前提是：SM上必须有足够的活跃warp。如果每个线程占用太多寄存器或共享内存，会导致单个SM能容纳的线程块数量减少，进而活跃warp不足，机器就会空转。
这就是“占用率”（occupancy）问题：追求单个线程效率（如用更多寄存器）可能牺牲整体并行度。GPU优化常陷入这种矛盾——你得在“每个线程干得漂亮”和“全场线程一起干”之间找平衡。有时故意降低单线程效率，反而能让整体吞吐翻倍。
所以性能优化的重点不是让单个线程飞快，而是要确保整个设备始终有足够工作。
分支分歧是隐形杀手，同束线程必须步调一致
GPU的线程执行单位是warp，32个线程必须同步执行同一条指令。
如果warp内线程走不同分支，硬件不得不分别跑不同分支代码，同时屏蔽不属于当前分支的线程，这样同样的逻辑会被重复执行。
如果代码中有if-else分支，而同一warp内的线程走了不同路径，灾难就来了：GPU会先执行if分支，禁用走else的线程；再执行else分支，禁用走if的线程。
虽然结果正确，但时间翻倍。在GPU代码里分支是很贵的，尤其当warp里线程不一致时。好的核函数设计尽量让warp内线程走同样路径，避免由于分支而浪费。这种“分支分歧”（divergence）在CPU上影响不大，但在GPU上足以让性能崩盘。
因此，高性能CUDA代码会刻意对齐数据，让相邻线程（尤其同一warp内）尽量走相同分支。比如处理图像时，按行或按块分配任务，避免随机跳跃导致控制流混乱。
内存访问要“排队”，散乱读写直接拖垮带宽
全局显存基于DRAM，最擅长连续大块读写。
全局内存带宽不仅取决于你读写多少数据，还取决于访问模式。DRAM喜欢按连续地址大块读取。当warp中的线程访问连续地址时，这些请求能被硬件合并成少量大块访问，效率极高；但如果访问很散乱，原本一条指令可能变成很多低效访问。
如果一个warp里32个线程各自读取分散的地址，GPU可能被迫发起32次独立内存请求，带宽利用率暴跌。反之，若线程按顺序读取连续地址（如thread 0读地址0，thread 1读地址4……），硬件会自动合并成1-2次高效传输，这就是“内存合并”（coalescing）。
很多算法天然访问模式不佳，但可通过共享内存中转：先由线程协作以合并方式把数据搬进共享内存，再按任意顺序读取。一旦数据上了片上高速存储，访问模式就不再受限制。这条规则决定了：即使算法逻辑不变，调整线程与数据的映射关系，也能带来数倍性能差异。
因此让线程访问连续内存地址，或者通过共享内存先组织好访问模式，是提升带宽利用的关键。
并行模型没有银弹，只有成本转移的艺术
从MPI到OpenMP再到CUDA，并行编程的本质不是技术竞赛，而是成本分配的选择。
MPI把成本压给程序员——手动管理节点间通信，换来超强扩展性；
OpenMP把成本压给硬件——靠缓存一致性协议简化编程，但难以扩展到上千核；
CUDA则把成本显式暴露给开发者：你必须操心数据搬运、内存层级、线程协作。
没有哪种模型绝对优越，关键看问题特性。如果你的任务能拆成百万独立子任务（如图像处理、矩阵运算），CUDA的成本付出换来巨大回报；如果任务高度串行或通信密集，硬上GPU反而得不偿失。理解这一点，才能跳出“该不该学CUDA”的纠结，转而思考“我的问题到底适合在哪里付成本”。
回归第一性原理：性能是设计出来的，不是等来的
这篇文章的价值，不在于教具体CUDA语法，而在于揭示现代计算性能的根本逻辑。当摩尔定律失效，硬件不再替你兜底，程序员必须重新理解机器本质。GPU的威力不在“快”，而在“多”；不在“智能调度”，而在“结构化并行”。
CUDA强迫你直面内存墙、带宽瓶颈、并行粒度等底层约束，从而写出真正高效代码。这种思维迁移，远比掌握某个API重要——无论未来硬件如何演进，理解“成本在哪、如何权衡”，永远是性能工程的核心。
最后的思考框架
回顾这些基础，会发现一个反复出现的主题：各种并行模型不是为了证明谁最优，而是在不同地方做权衡，是在硬件、软件、人力投入之间选择付出成本。CUDA之所以让设计变得清晰，是它强迫你关注这些权衡点：数据在哪、何时搬、如何并行组织、如何提升数据重用。
所以学CUDA的核心不是为了学一个API，而是为了理解这个架构和如何评估你问题本身的限制条件，并在这些条件下选择合适付出。