一些前置的东西

SIMD(Single Instruction, Multiple Data)

  • 属于硬件/指令集层面，是硬件能力，一条指令对一排数据做同样的计算。要求所有数据走完全一致的运算路径，不能各走各的分支

SIMT(Single Instruction, Multiple Threads)

  • 是NVIDIA为GPU造的词，属于硬件执行模型层，是硬件和编程模型中间的“桥”。一个Wrap(32个线程)共用一条指令流，但每个线程有自己的寄存器，自己的内存地址，甚至可以走不同分支。本质是用SIMD的底层硬件，封装成一堆看起来独立的线程。

SPMD(Single Program, Multiple Data)

  • 是变成模型，和硬件解耦。所有进程/线程跑同一份代码，处理不同数据。运行时可以走各自的不同路径，停在不同位置。这不要求指令级同步。它可以跑在完全独立的多核/集群硬件上

  • 写CUDA使用SPMD思维，GPU使用SIMT执行模型调度，硬件底层通过SIMD式的向量单元。三层串成一条链

Flynn 分类法

  • 这种分类法将 “数据流 x 指令流”将架构分为四类
    • SISD: 传统单核顺序执行
    • SIMD： 单指令多数据，如上
    • MISD： 多指令处理同一条数据，及其罕见
    • MIMD： 每个核跑自己的指令，自己的数据，多核CPU，计算集群多是它
  • 在编程模型层，还有：MPMD：不同节点跑不同程序，比如master-worker架构，或者耦合仿真
  • 其他类似的东西：
    • SWAR(SIMD within a Register)：不用专门的向量指令，将多个小数据塞进一个普通通用寄存器手动并行，是“穷人版”的SIMD
    • SMT(Simultaneous Multithreading)：一个物理核同时跑多个线程来填满流水线空袭，是线程并行级技术，不在Flynn分类中

并行范式

从底层到顶层分别是：
体系结构 -> 变成模型 -> 业务场景

层级分类 分类标准 对应内容
体系结构 硬件原生的并行能力，并行的物理根源 SIMD、 MIMD (SPMD)
编程模型 开发者视角的通用并行逻辑，硬件与软件的桥梁 数据并行、 任务并行、 线程并行
务场景 针对特定业务/计算场景的专属并行模式 管道并行、MapReduce、图计算并行
  • GPU适合什么样的并行范式？
  • GPU的设计哲学：多小核，弱控制，弱缓存，强并行。
    GPU本质上是SIMT架构，这一特性决定了计算范式：
    • 同质化
    • 高并行
    • 无复杂分支
    • 规整内存访问

图形渲染并行是GPU的原生设计目标，这一点是显然的，不然为什么要叫GPU
* 顶点处理，片元处理天然具备百万级同质化数据并行度，每个顶点/像素处理逻辑完全一致，与GPU的SIMT架构匹配

数据并行(GPGPU核心范式)

数据并行为GPGPU时代的核心范式，是GPU从图形专用扩展至通用计算的核心实现形式

SIMD/SPMD

这是GPU硬件底层的原生执行范式， 是所有高层并行范式的硬件基础

管道并行/流并行

管道并行是GPU流式硬件架构原生支持的并行范式

GPU不适合的并行范式

  • 细粒度任务并行，分支密集型，通用线程级并行，都属于GPU的短板
    这一类范式有复杂的控制逻辑，大量分支判断，差异化任务执行逻辑。这些导致Warp内线程分化，造成大量计算核心限制，导致硬件资源利用率低，执行性能低于单核CPU

GPU

GPU是为了高吞吐量而优化的，和CPU的设计相反。GPU擅长同时处理海量任务
一颗完整的GPU含有以下模块，通过片上互联总线相连

模块 功能
Streaming Multiprocessor (SM) GPU的计算核心单元，总算力=Σ(SM)
内存控制器 负责连接显存与芯片的读写
全局调度器 将用户提交的计算任务动态分配给SM执行
层次化存储体系 显存、全局内存等多层次存储结构

SM是GPU调度的最小单元，也是GPU物理架构的核心基本快，每一个SM是一个完整的计算单元，包含以下核心组件：

组件 功能
Warp调度器+指令分发单元 将线程块拆分为Warp，调度Warp执行
StreamProcessor(SP) GPU最小物理执行单元，极简ALU
专用加速单元 SFU等
寄存器文件 SM内最快存储，给活跃线程提供私有存
L1缓存／共享内存 可配置片上存储，灵活分配
加载/存储单元 处理线程的内存读写请求

硬件执行单元Warp

  • warp是GPU硬件调度和执行的最小单元，由硬件自动生成，固定包含32个Thread
    GPU采用SIMT架构，一个Warp内所有Thread永远同步执行同一条指令，但是每个Thread处理不同的数据

  • 注意：分支发散：

    • 即使代码中存在分支，warp也会执行所有分支导致空转，从而影响性能
      当定义一个包含N个Thread的Block时，GPU会自动将Block内的Thread按顺序每32个组成一个Warp。因此block中的thread总数一般设置为32的证书被，避免Warp资源浪费

CUDA

CUDA是NVIDIA推出的并行计算平台与编程模型，以C++为基础进行语言扩展，同时提供了一套完整的API，用于管理GPU设备

CUDA的软件线程模型层级

CUDA在软件层面定义了一套清晰的层级化线程结构，与硬件执行单元解耦，这套层级结构是：
Grid -> Block -> Thread

Thread

Thread 是 kernel函数的一次执行实例，也是CUDA中最小的逻辑执行单元。
- 同一kernel内所有Thread都会执行完全相同的代码，但是Thread通过自己唯一ID，处理数据集中的不同部分
- 逻辑上Thread是最小执行单位，但是GPU永远不会单独调度一个Thread

Block

Block是一组Thread组成的中间层逻辑单元，也是GPU将任务分配给SM的最小软件单位
- 同一个block内的所有thread会被分配到同一个SM上执行
- BLock内的thread可通过SM内的共享内存实现高速数据共享，也可以通过 __syncthreads() 等进行线程间同步
- block支持 1D/2D/3D但是block内tread总数有硬件上限

Grid

Grid是 GPU执行一个kernel函数的最上层逻辑单元，一个kernel函数唯一对应一个Grid
- 同一个Grid的Block，都执行同一个Kernel函数，共同处理同一个计算任务
- Grid同样支持1D/2D/3D索引结构
- Grid内所有Block完全独立，并行执行，无固定顺序。不同Block之间只能通过全局显存通信，无法高效共享数据