CUDA学习之路3——GPU并行本质与软硬件核心逻辑

3962 字
20 分钟
CUDA学习之路3——GPU并行本质与软硬件核心逻辑

引言#

在前两篇内容中1.速通环境配置2.C++关键知识点复习
我们尝试搭建了CUDA的开发环境,并且复习了C/C++的一些核心概念。那么很多同学实操后会产生一些疑问:

  1. 为什么这段代码能在 GPU 上并行执行?
  2. GPU 和 CPU 的本质区别到底是什么?
  3. CUDA 究竟扮演了怎样的角色?
  4. 究竟什么时候用GPU来计算?

这篇文章,我们从GPU的硬件设计本源出发,厘清并行计算的完整分层逻辑,拆解GPU和CUDA的全栈体系,帮大家建立
硬件架构 - 编程模型 - 业务场景」的完整认知,为后续编写高性能CUDA代码打下最坚实的基础。

并行计算#

并行范式#

那么哪一类计算是能够被并行化的呢?

诸如SIMD、SIMT、等等令人头疼的名字经常令我们陷入混乱:哪些是底层硬件能力?哪些是编程模型?哪些又是场景封装?

我们先遵循一个逻辑,从底层到顶层分别是:

体系结构➡️编程模型➡️业务场景
层级分类分类标准对应内容
体系结构根据体系结构的指令与数据流来提供硬件原生的支持的并行能力,是所有并行的物理根源单指令多数据流SIMD、多指令多数据流MIMD(单程序多数据流SPMD)
编程模型开发者视角的通用并行逻辑,是硬件与软件之间的桥梁,不会绑定特定的业务数据并行、任务并行、线程并行
业务场景针对特定的业务/计算场景的专属并行模式,是基于特定业务的场景化封装管道并行、MapReduce并行、图计算并行

经典的弗林分类法中的SISD为纯串行、MISD无应用场景,而SPMD是MIMD最主流的核心范式,因此纳入。

那么从GPU来看,它究竟适合什么样子的并行范式呢?

这些范式又是如何跟着GPU架构的演进而发展呢?

我们提过GPU的设计思路是多小核、弱控制、弱缓存、强并行。 那么实际上,GPU的从硬件来看是SIMT(单指令多线程),是SIMD的变体。

这个硬件决定了它的计算范式:同质化、高并行、无复杂分支、规整内存访问的并行范式。

哪些范式适合GPU呢?#

适合GPU的并行范式
适合GPU的并行范式

1. 图形渲染并行#

图形渲染并行是GPU的原生设计目标,也是所有GPU并行范式的技术起源。

技术衍伸

GPU的诞生,核心功能定位就是图形渲染。1990年代出现的固定功能渲染管线,已实现硬件级的管道并行与SIMD范式;后续技术演进中,从固定管线到可编程着色器,再到当前的硬件光线追踪,图形渲染并行始终是GPU的原生核心能力,也是后续所有通用并行范式的技术起点。

适配逻辑

图形渲染的顶点处理、片元处理环节,天然具备百万级的同质化数据并行度,每个顶点、每个像素的处理逻辑完全一致,与GPU的SIMT架构完全匹配;渲染管线的阶段拆分,属于硬件级的管道并行,调度逻辑完全由GPU硬件实现,无额外软件调度开销。

2. 数据并行#

数据并行是GPGPU(通用GPU计算)时代的核心范式,是GPU从图形专用芯片扩展至通用计算领域的核心实现形式。

技术衍伸

2007年NVIDIA发布CUDA架构,正式开启GPGPU时代。CUDA的核心设计,是将GPU图形渲染中成熟应用的SIMD并行能力,抽象为通用的数据并行编程模型,使开发者无需掌握图形API,即可通过类C语言编写GPU并行代码。深度学习领域的技术爆发,完全建立在GPU数据并行能力的基础上,神经网络的张量计算,属于极致的大规模数据并行。

适配逻辑

数据并行的核心执行逻辑为同一执行逻辑,分片处理数据,与GPU的SIMT架构完全匹配。CUDA的核函数编程,是典型的SPMD数据并行实现。

主流场景

深度学习训练与推理、矩阵运算、图像处理、数值模拟,是当前GPU使用最广泛的并行范式。

3. SIMD/SPMD#

SIMD/SPMD是GPU硬件底层的原生执行范式,是上述所有高层并行范式的硬件基础。

技术衍伸

SIMD是GPU硬件架构的底层核心,从最早的2D图形加速卡,到当前的H200、A800等高端计算卡,GPU的硬件计算核心始终基于SIMD架构构建。SPMD是CUDA编程模型的核心范式。

适配逻辑

GPU的向量指令集、张量核心,均基于SIMD架构设计,针对FP16/BF16精度的张量运算,单条指令即可完成数百个乘加运算,计算密度远超CPU。SPMD编程模型,与GPU的分布式SM架构完全适配,可扩展性极强,从单卡到多卡集群,均可复用同一套SPMD执行逻辑。

4. 管道并行/流并行#

管道并行(也叫流并行),是GPU流式硬件架构原生支持的并行范式。

技术衍伸

管道并行最早是图形渲染管线的原生执行模式,随着深度学习技术的发展,管道并行被扩展至多卡GPU训练场景,用于解决大模型训练的显存瓶颈;同时在视频流处理、实时推理场景,管道并行也是GPU的常用执行范式。

适配逻辑

GPU的流式多处理器、多流调度机制,原生支持管道并行。多个计算流可在不同的SM上,并行执行流水线的不同阶段,通过事件同步机制即可实现无额外软件开销的流水线调度。在大模型训练场景中,管道并行将模型的不同层拆分到不同GPU上,实现流水线式的前向传播与反向传播,可同时解决大模型训练的显存限制与执行效率问题。

5. GPU完全不适合的并行范式#

该模块内的并行范式与GPU的硬件架构设计完全相悖,不仅无法获得性能提升,甚至可能出现执行性能远低于单核CPU的情况,是CPU多核架构的核心优势领域。

细粒度任务并行、分支密集型、通用线程级并行,均属于GPU的绝对短板。 这类范式的核心特征,是复杂的控制逻辑、大量的分支判断、差异化的任务执行逻辑。这些特征会直接导致GPU的Warp内线程分化,进而造成大量计算核心闲置,硬件资源利用率极低,最终执行性能甚至不如单核CPU。


GPU是什么?#

Graphics Processing Unit (GPU) 图形处理器,是专为图形渲染、大规模并行计算场景设计的微处理器,如今已成为 AI 训练、科学计算、高性能计算领域的核心算力载体。

GPU在计算体系的定位#

GPU 与 CPU 的核心差异,根植于完全相反的设计哲学。

CPU和GPU的设计差异
CPU和GPU的设计差异

Nvidia的首席科学家Bill Dally说过:

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从物理芯片层面拆解了GA102的晶体管与模块,而Intro to GPUs更注重逻辑抽象,建立更加清晰的执行模型。

一颗完整的GPU芯片,物理上由以下模块组成,并且这些模块通过片上互联总线相连:

模块功能
Streaming MultiProcessor (SM)GPU的计算核心单元,互相之间独立,GPU的总算力=SUM(SM)
内存控制器负责连接显存与芯片的读写
全局调度器GPU的控制器,负责把用户提交的计算任务动态分配给SM执行
层次化存储体系包括显存,全局内存等多层次的存储结构

Streaming Multiprocessor
Streaming Multiprocessor

SM 是GPU 硬件调度任务的最小单元,也是 GPU 物理架构的核心基本块。NVIDIA 官方命名为 SM,AMD 对应叫计算单元 CU(Compute Unit),二者设计理念完全一致,行业通用 SM 代指这一核心模块。

每个SM都是一个完整的计算单元,物理内部包含以下核心组件:

组件功能
Warp调度器+指令分发单元SM的调度器,负责把分配到SM的线程块拆分成为Warp,调度Warp的执行
Stream Processor (SP)GPU的最小物理执行单元,是一个极简的算术逻辑单元ALU
专用加速单元SFU
寄存器文件SM 内速度最快的存储,容量可达数百 KB,给活跃线程提供私有存储
L1 缓存 / 共享内存SM 内的可配置片上存储,可灵活分配为 L1 数据缓存或可编程共享内存
加载/存储单元处理线程的内存读写请求

硬件执行单元Warp#

它是 GPU 硬件调度和执行的最小物理单元,由硬件自动生成,固定包含32个Thread。

GPU采用SIMT,一个Warp内的所有Thread永远都同步执行同一条指令,但是每个Thread处理不同的数据。

注意

即使代码中存在分支(如 if-else),Warp 内的 Thread 也会执行所有分支(只是不满足条件的 Thread 会 “空转”,不写回结果),这就是分支发散,会影响性能。

当程序员定义了一个包含N个Thread的Block时,GPU硬件会自动将Block内的Thread按顺序每32个组成一个Warp。

例如Block有100个Thread,硬件就会生成4个Warp,有3个是满的,剩下的一个不满,只有4个活跃的Thread。不满的这个仍然会占用一个完整的Warp资源,这就导致了算力浪费。

因此写代码时一般建议将Block的Thread总数设置为32的整数倍。
SM的示意图
SM的示意图
Hierarchy of theads running on a GPU
Hierarchy of theads running on a GPU

CUDA是什么?#

CUDA是NVIDIA推出的并行计算平台与编程模型,它以C++为基础进行语言扩展,同时提供了一套完整的API,用于管理GPU设备(包括显存分配、内核函数启动、设备同步等)。

只要应用场景包含可并行处理的数据,CUDA就能通过大规模并发线程显著提升计算性能。

CUDA的软件线程模型层级#

为了组织大规模并行任务,CUDA 在软件层面定义了一套清晰的层级化线程结构。这是程序员在代码中直接控制的逻辑模型,与硬件执行单元解耦。

这套层级结构是:Grid➡️Block➡️Thread。

1. Thread#

Thread 是 Kernel 函数的一次执行实例,也是 CUDA 中最小的逻辑执行单元。

同一 Kernel 内的所有 Thread,都会执行完全相同的代码(SPMD,Single Program Multiple Data 思想),但每个 Thread 通过自己的唯一 ID(threadIdx)处理数据集中的不同部分。
例如:给长度为 10000 的数组每个元素 +1,可启动 10000 个 Thread,每个 Thread 仅处理数组中的一个元素。
注意

逻辑上Thread是最小执行单位,但是硬件GPU永远都不会单独调度单个Thread。

2. Block#

Block 是由一组 Thread 组成的中间层逻辑单元,也是 GPU 将任务分配给 SM(流式多处理器,Streaming Multiprocessor)的最小软件单位。

同一 Block 内的所有 Thread,会被分配到同一个 SM 上执行。
Block 内的 Thread 可通过 SM 内的共享内存实现高速数据共享,还可通过 __syncthreads()等原语进行线程间同步,这是 GPU 并行算法的核心基础。
限制

Block 支持 1D/2D/3D 的索引结构,但 Block 内的 Thread 总数有硬件上限(常见为 1024,具体取决于 GPU 架构)。

3. Grid#

Grid 是 GPU 执行一个 Kernel 函数的最上层逻辑单元,一个 Kernel 函数唯一对应一个 Grid。

Grid 由一组 Block 组成,所有属于同一 Grid 的 Block,都执行同一个 Kernel 函数,共同处理同一个计算任务。
Grid 同样支持 1D/2D/3D 的索引结构,且与 Block 的维度对应。例如处理 1920×1080 的图片,可用 120×68 个 2D Block 组成一个 2D Grid。
注意

Grid 内的所有 Block 完全独立、并行执行,无固定顺序。GPU 的全局调度器会动态将 Block 分配给空闲的 SM,不同 Block 之间只能通过全局显存通信,无法高效共享数据。


后言#

经过上述软硬件学习,我们大概了解了一些最最基本的CUDA和GPU的概念,那么下一节我们就正式开始CUDA编程学习!

支持与分享

如果这篇文章对你有帮助,欢迎分享给更多人或赞助支持!

赞助
CUDA学习之路3——GPU并行本质与软硬件核心逻辑
https://dlog.com.cn/posts/cuda03/cuda编程概述/
作者
杜子源
发布于
2026-03-27
许可协议
CC BY-NC-SA 4.0
Profile Image of the Author
杜子源
都是风景,幸会
公告
如果需要原图,请私信联系我。
音乐
封面

音乐

暂未播放

0:00 0:00
暂无歌词
分类
标签
站点统计
文章
6
分类
3
标签
4
总字数
11,879
运行时长
0
最后活动
0 天前

目录