CUDA学习之路[3]——GPU并行本质 | 硬件架构 × 编程模型
目录
引言
在前几篇内容中我们尝试搭建了CUDA的开发环境，并且复习了C/C++的一些核心概念。
很多同学实操后会产生一些疑问：
- 为什么这段代码能在 GPU 上并行执行？
- GPU 和 CPU 的本质区别到底是什么？
- CUDA 究竟扮演了怎样的角色？
- 究竟什么时候用GPU来计算？
- 理解并行计算的分层逻辑（体系结构 → 编程模型 → 业务场景）
- 厘清GPU的硬件架构原理（SM、Warp、SP）
- 掌握CUDA的软件线程模型（Grid、Block、Thread）
本文从GPU的硬件设计本源出发，为后续编写高性能CUDA代码打下基础。
并行计算：你需要在哪一层思考并行？
并行范式
那么哪一类计算是能够被并行化的呢？
SIMD、SIMT、SPMD……这些令人头疼的名字经常让我们陷入混乱：哪些是底层硬件能力？哪些是编程模型？哪些又是场景封装？
我们先遵循一个逻辑，从底层到顶层分别是：
| 层级分类 | 分类标准 | 对应内容 |
|---|---|---|
| 体系结构 | 硬件原生的并行能力，并行的物理根源 | SIMD、MIMD（SPMD） |
| 编程模型 | 开发者视角的通用并行逻辑，硬件与软件的桥梁 | 数据并行、任务并行、线程并行 |
| 业务场景 | 针对特定业务/计算场景的专属并行模式 | 管道并行、MapReduce、图计算并行 |
经典的弗林分类法中，SISD为纯串行、MISD无实际应用场景，而SPMD是MIMD最主流的核心范式。
GPU适合什么样的并行范式？
GPU的设计哲学是 多小核、弱控制、弱缓存、强并行。
从硬件层面看，GPU本质上是 SIMT（单指令多线程） 架构，这是SIMD的变体。这一硬件特性决定了它的计算范式： 同质化、高并行、无复杂分支、规整内存访问
哪些范式适合GPU？
1. 图形渲染并行（原生范式）
图形渲染并行是GPU的原生设计目标，也是所有GPU并行范式的技术起点。
GPU的诞生核心功能就是图形渲染。1990年代的固定功能渲染管线已实现硬件级的管道并行与SIMD范式；从固定管线到可编程着色器，再到当前的硬件光线追踪，图形渲染并行始终是GPU的原生核心能力。
顶点处理、片元处理环节天然具备百万级的同质化数据并行度，每个顶点/像素的处理逻辑完全一致，与GPU的SIMT架构完美匹配。
2. 数据并行（GPGPU核心范式）
数据并行是GPGPU时代的核心范式，是GPU从图形专用芯片扩展至通用计算领域的核心实现形式。
2007年NVIDIA发布CUDA架构，正式开启GPGPU时代。CUDA将GPU图形渲染中的SIMD并行能力，抽象为通用的数据并行编程模型，使开发者无需掌握图形API即可编写GPU并行代码。
数据并行的核心执行逻辑为 同一执行逻辑，分片处理数据，与GPU的SIMT架构完全匹配。CUDA的核函数编程是典型的SPMD数据并行实现。
深度学习训练与推理、矩阵运算、图像处理、数值模拟。
3. SIMD/SPMD（硬件底层范式）
这是GPU硬件底层的原生执行范式，是所有高层并行范式的硬件基础。
SIMD是GPU硬件架构的底层核心，从2D图形加速卡到当前的H200、A800，GPU的计算核心始终基于SIMD构建。SPMD则是CUDA编程模型的核心范式。
4. 管道并行 / 流并行
管道并行是GPU流式硬件架构原生支持的并行范式。
管道并行最早是图形渲染管线的原生模式，后被扩展至多卡GPU训练场景，用于解决大模型训练的显存瓶颈。
GPU的流式多处理器、多流调度机制原生支持管道并行。在大模型训练中，管道并行将模型不同层拆分到不同GPU上，实现流水线式的前向与反向传播。
5. GPU完全不适合的并行范式
细粒度任务并行、分支密集型、通用线程级并行，均属于GPU的绝对短板。
这类范式的核心特征是复杂的控制逻辑、大量的分支判断、差异化的任务执行逻辑。这些特征会直接导致Warp内线程分化，造成大量计算核心闲置，硬件资源利用率极低，执行性能甚至不如单核CPU。
GPU是什么？
Graphics Processing Unit (GPU) 图形处理器，是专为图形渲染、大规模并行计算场景设计的微处理器，如今已成为AI训练、科学计算、高性能计算领域的核心算力载体。
GPU在计算体系中的定位
GPU与CPU的核心差异，根植于完全相反的设计哲学。
CPU是为了低延迟优化的，以获得非常棒的单线程性能；而GPU是为了高吞吐量优化的，擅长同时处理海量任务。
CPU ≈ 将军：啥都能干，但个数有限，纵使齐天之资也分身乏力
GPU ≈ 千军万马：只会标准化的傻瓜操作，但有千军万马，一人一滴血都能耗死将军。
那么为了让大家更清晰的理解GPU在整个计算芯片中的定位，我整理了主流的芯片进行对比：
| 芯片类型 | 设计哲学 | 优势&劣势 | 典型应用场景 | 代表产品/厂商 |
|---|---|---|---|---|
| CPU | 通才 逻辑控制和串行计算，低延迟 | 通用性强，生态成熟，单核响应快 并行算力弱，能效比低 | 操作系统、数据库、任务调度 | Intel Xeon, AMD EPYC |
| GPU | 将才 大规模并行 | 并行计算能力极强，生态完善 功耗高，单核效率不如CPU | AI模型训练、科学计算、图形渲染 | NVIDIA H100, AMD MI300X |
| FPGA | 塑才 硬件可重构，电路仿真 | 灵活性高，延迟极低，能效优于GPU 开发难度大，需硬件语言，峰值算力较低 | 算法未定型、通信基站、高频交易 | Xilinx (AMD) Versal, Intel Stratix |
| ASIC | 专才 为特定任务全定制设计 | 性能、能效、成本（量产时）达到极致 前期研发成本极高，流片后功能无法修改 | 算法成熟、大规模量产场景（如挖矿） | 各类定制芯片 |
| TPU | 怪才 谷歌的专用ASIC，专为NN中的张量运算设计 | 针对AI计算的能效比极高，吞吐量大 生态相对封闭，主要绑定Google Cloud | 大规模云端AI推理与训练 | Google TPU v4/v5 |
| NPU | 辅才 专注于边缘端AI推理，通常集成在手机/汽车SoC中 | 极低功耗，高能效，响应快，保护数据隐私 算力相对有限，主要用于推理而非训练 | 手机人脸识别、语音助手、自动驾驶 | 华为昇腾, 苹果A17 Pro, 高通Hexagon |
| PIM | 隐才 解决存储墙问题 | 极大幅度减少数据搬运，延迟极低，功耗极低 技术前沿，软件生态不成熟 | 内存带宽瓶颈严重的场景，AI推理 | SK海力士AiM, 三星HBM-PIM |
GPU的核心架构
结合两篇高质量资料：Inside the GPU 从物理芯片层面拆解，Intro to GPUs 更注重逻辑抽象。
一颗完整的GPU芯片，物理上由以下模块组成，通过片上互联总线相连：
| 模块 | 功能 |
|---|---|
| Streaming Multiprocessor (SM) | GPU的计算核心单元，总算力 = Σ(SM) |
| 内存控制器 | 负责连接显存与芯片的读写 |
| 全局调度器 | 将用户提交的计算任务动态分配给SM执行 |
| 层次化存储体系 | 显存、全局内存等多层次存储结构 |
SM 是GPU 硬件调度任务的最小单元，也是 GPU 物理架构的核心基本块。NVIDIA 官方命名为 SM，AMD 对应叫计算单元 CU（Compute Unit），二者设计理念完全一致，行业通用 SM 代指这一核心模块。
每个SM都是一个完整的计算单元，物理内部包含以下核心组件：
| 组件 | 功能 |
|---|---|
| Warp调度器 + 指令分发单元 | 将线程块拆分为Warp，调度Warp执行 |
| Stream Processor (SP) | GPU最小物理执行单元，极简ALU |
| 专用加速单元 | SFU等 |
| 寄存器文件 | SM内最快存储，给活跃线程提供私有存储 |
| L1缓存 / 共享内存 | 可配置片上存储，灵活分配 |
| 加载/存储单元 | 处理线程的内存读写请求 |
硬件执行单元 Warp
Warp 是GPU硬件调度和执行的最小物理单元，由硬件自动生成，固定包含32个Thread。
GPU采用SIMT架构，一个Warp内的所有Thread永远同步执行同一条指令，但每个Thread处理不同的数据。
即使代码中存在分支（如if-else），Warp内的Thread也会执行所有分支（不满足条件的Thread会“空转”），这就是分支发散，会影响性能。
当程序员定义了一个包含N个Thread的Block时，GPU硬件会自动将Block内的Thread按顺序每32个组成一个Warp。
写代码时一般建议将Block的Thread总数设置为32的整数倍，避免Warp资源浪费。
CUDA是什么？
CUDA 是NVIDIA推出的并行计算平台与编程模型，以C++为基础进行语言扩展，同时提供了一套完整的API，用于管理GPU设备（显存分配、内核函数启动、设备同步等）。
只要应用场景包含可并行处理的数据，CUDA就能通过大规模并发线程显著提升计算性能。
CUDA的软件线程模型层级
为了组织大规模并行任务，CUDA在软件层面定义了一套清晰的层级化线程结构。这是程序员在代码中直接控制的逻辑模型，与硬件执行单元解耦。
这套层级结构是：Grid → Block → Thread
1. Thread
Thread 是Kernel函数的一次执行实例，也是CUDA中最小的逻辑执行单元。
同一Kernel内的所有Thread，都会执行完全相同的代码，但每个Thread通过自己的唯一ID（threadIdx）处理数据集中的不同部分。
给长度为10000的数组每个元素+1，可启动10000个Thread，每个Thread仅处理数组中的一个元素。
逻辑上Thread是最小执行单位，但硬件GPU永远都不会单独调度单个Thread。
2. Block
Block 是由一组Thread组成的中间层逻辑单元，也是GPU将任务分配给SM的最小软件单位。
同一Block内的所有Thread，会被分配到同一个SM上执行。
Block内的Thread可通过SM内的共享内存实现高速数据共享，还可通过__syncthreads()等原语进行线程间同步。
Block支持1D/2D/3D索引结构，但Block内的Thread总数有硬件上限（常见为1024，取决于GPU架构）。
3. Grid
Grid 是GPU执行一个Kernel函数的最上层逻辑单元，一个Kernel函数唯一对应一个Grid。
所有属于同一Grid的Block，都执行同一个Kernel函数，共同处理同一个计算任务。
Grid同样支持1D/2D/3D索引结构。例如处理1920×1080的图片，可用120×68个2D Block组成一个2D Grid。
Grid内的所有Block完全独立、并行执行，无固定顺序。不同Block之间只能通过全局显存通信，无法高效共享数据。
小结
通过本文，我们建立了从硬件架构（SM、Warp、SP） 到编程模型（Grid、Block、Thread） 的完整认知链路。
掌握了GPU的并行本质和CUDA的线程模型后，下一节我们将正式进入CUDA编程实战，从第一个Kernel函数开始，逐步深入内存管理、性能优化等核心主题。
支持与分享
如果这篇文章对你有帮助，欢迎分享给更多人或赞助支持！
部分内容可能已过时
音乐
暂未播放