登录注册
京东自营 + 国补 iPhone 历史最低价          国家补贴 享8折

GPU到底是如何工作的？这篇AI Infra入门全部告诉你

欢马劈雪

工程师 (已认证)

  • 原创分享签约作者

发表于教程文章
102

背景

AI 流行的当下，你有没有想过：

大模型推理服务到底怎么跑起来的？大模型推理服务的运行过程中，CPU和GPU分别负责哪些工作？
用GPU一定比CPU跑的快么？哪些场景需要用GPU?

图形渲染到GPGPU

为图形而生

GPU最初的使命是加速图形渲染。而渲染一帧图像，本质上就是对数百万个像素点进行相似的计算，这天然就是一种大规模并行任务。

可编程性的开启 (2001)

NVIDIA发布GeForce 3，首次引入可编程着色器 (Programmable Shaders)。实质上允许开发者为 GPU 编写软件，让GPU的众多并行处理单元去同时执行，以精确控制光照和颜色如何加载到显示器上。这是朝着加速计算方向迈出的重要一步，因为它允许开发者直接为 GPU 编写软件。

学术界的探索

一批敏锐的研究人员意识到，GPU的本质就是一个拥有数百甚至数千个核心的大规模并行架构，其浮点运算吞吐量远超当时的CPU。他们的核心想法是：能不能用GPU进行科学计算? 开始探索利用GPU计算科学计算问题，从而利用GPU的算力。这便是GPGPU（通用计算GPU）的萌芽。但是门槛非常高， 需要开发者同时精通图形学和科学计算。

NVIDIA的抉择

NVIDIA敏锐地捕捉到了GPGPU的发展潜力，开始不再局限于加速图形渲染，主动拥抱GPGPU。

2006年，发布了第一款为通用计算设计的统一架构GPU - GeForce 8800 GTX显卡（G80架构）。它将GPU内部的计算单元统一起来，形成了一个庞大的、灵活的并行核心阵列，为通用计算铺平了硬件道路。

2007年，NVIDIA正式推出了CUDA平台。CUDA的革命性在于，它提供了一套简单的编程模型，让开发者能用近似C语言的方式，轻松地驾驭GPU内部成百上千个并行核心。 开发者无需再关心复杂的图形接口，可以直接编写在数千个线程上并发执行的程序。至此终结了GPGPU编程的蛮荒时代，让GPU计算真正走下神坛，成为开发者触手可及的强大工具。

随着深度学习的发展与流行，CUDA生态系统目前也成为NVIDIA最深、最宽的护城河。

参考链接 nvidia-past-present-and-future

CPU/GPU异构计算架构

CPU是整个系统的核心，是总指挥，GPU的任务指令是由CPU分配的。

CPU通过PCIe总线给GPU发送指令和数据交互。而PCIe支持DMA和MMIO两种通讯模式:

  • MMIO（内存映射I/O）由CPU直接控制数据读写，操作系统会把设备地址映射到CPU的虚拟空间中，适合小数据量的指令交互
  • DMA（直接内存访问）则允许设备绕过CPU直接访问系统内存，专为大数据块的高效传输设计。

CPU通过IMC和Memory Channel访问内存，为了提升数据传输带宽，高端CPU通常会支持多内存通道，即多IMC和Memory Channel的组合，以满足日益增长的数据处理需求。

一个简单的应用

讲道理，对于开发来说，再通俗易懂的语言描述都不如一个简单Demo来的实在。

Demo代码来自even-easier-introduction-cuda,可在collab测试运行下述代码。

实现两个长度为 23? (约10亿) 的浮点数数组的相加。其中，一个数组 (x) 的所有元素初始化为 1.0，另一个数组 (y) 的所有元素初始化为 2.0，我们计算 y[i] = x[i] + y[i]。

CPU的实现

#include <iostream>
#include <math.h>
#include <chrono>

// function to add the elements of two arrays
void add(int n, float *x, float *y)
{
  for (int i = 0; i < n; i++)
      y[i] = x[i] + y[i];
}

int main(void)
{
  int N = 1<<30;

  float *x = new float[N];
  float *y = new float[N];

  // initialize x and y arrays on the host
  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  auto start = std::chrono::high_resolution_clock::now();

  // Run kernel on 1M elements on the CPU
  add(N, x, y);

  auto stop = std::chrono::high_resolution_clock::now();
  auto duration = std::chrono::duration_cast<std::chrono::milliseconds>(stop - start);
  std::cout << "CPU 'add' function execution time: " << duration.count() << " ms" << std::endl;

  // Check for errors (all values should be 3.0f)
  float maxError = 0.0f;
  for (int i = 0; i < N; i++)
    maxError = fmax(maxError, fabs(y[i]-3.0f));

  std::cout << "Max error: " << maxError << std::endl;

  delete [] x;
  delete [] y;

  return 0;
}

性能表现

g++ add.cpp -o add
time ./add

CPU 'add' function execution time: 3740 ms
Max error: 0

real 0m21.418s
user 0m15.798s
sys 0m4.400s
  • 计算耗时: 核心的add函数耗时 3740毫秒。
  • 总耗时: 整个程序从启动到结束（real time）耗时 21.4秒。这额外的时间主要消耗在分配8GB内存（new float[N]）以及初始化数组上。

GPU的实现

这里的代码后面会详细解读，此处看懂含义即可。

  1. 分配内存: 分别在CPU（Host）和GPU（Device, cudaMalloc）上分配内存。
  2. 数据传输 (H2D): 将CPU上的输入数据 (h_x, h_y) 拷贝到GPU显存 (d_x, d_y)。
  3. 执行Kernel函数: 在GPU上启动addKernel函数，利用其大规模并行能力进行计算。
  4. 数据传回 (D2H): 将GPU计算完成的结果 (d_y) 拷贝回CPU内存 (h_y) 以便后续使用或验证。
#include <iostream>
#include <math.h>

#define CUDA_CHECK(call) \
do { \
    cudaError_t err = call; \
    if (err != cudaSuccess) { \
        fprintf(stderr, "CUDA Error in %s at line %d: %s\n", __FILE__, __LINE__, cudaGetErrorString(err)); \
        exit(EXIT_FAILURE); \
    } \
} while (0)


// __global__ 关键字声明的函数被称为Kernel函数
__global__
void add(int n, float *x, float *y)
{
  int index = blockIdx.x * blockDim.x + threadIdx.x;
  if (index < n) {
    y[index] = x[index] + y[index];
  }
}

int main(void)
{
  int N = 1 << 30;
  size_t bytes = N * sizeof(float);


  float *h_x, *h_y;
  h_x = new float[N];
  h_y = new float[N];

  float *d_x, *d_y;
  CUDA_CHECK(cudaMalloc(&d_x, bytes));
  CUDA_CHECK(cudaMalloc(&d_y, bytes));

  for (int i = 0; i < N; i++) {
    h_x[i] = 1.0f;
    h_y[i] = 2.0f;
  }


  CUDA_CHECK(cudaMemcpy(d_x, h_x, bytes, cudaMemcpyHostToDevice));
  CUDA_CHECK(cudaMemcpy(d_y, h_y, bytes, cudaMemcpyHostToDevice));

  cudaEvent_t start, stop;
  CUDA_CHECK(cudaEventCreate(&start));
  CUDA_CHECK(cudaEventCreate(&stop));
  CUDA_CHECK(cudaEventRecord(start));

  int blockSize = 256;
  int numBlocks = (N + blockSize - 1) / blockSize;
  add<<<numBlocks, blockSize>>>(N, d_x, d_y);

  CUDA_CHECK(cudaEventRecord(stop));
  CUDA_CHECK(cudaEventSynchronize(stop));

  float milliseconds = 0;
  CUDA_CHECK(cudaEventElapsedTime(&milliseconds, start, stop));
  std::cout << "GPU Kernel 'add' execution time: " << milliseconds << " ms" << std::endl;


  CUDA_CHECK(cudaEventDestroy(start));
  CUDA_CHECK(cudaEventDestroy(stop));
  CUDA_CHECK(cudaMemcpy(h_y, d_y, bytes, cudaMemcpyDeviceToHost));

  float maxError = 0.0f;
  for (int i = 0; i < N; i++) {
    maxError = fmax(maxError, fabs(h_y[i] - 3.0f));
  }
  std::cout << "Max error: " << maxError << std::endl;

  delete[] h_x;
  delete[] h_y;
  
  CUDA_CHECK(cudaFree(d_x));
  CUDA_CHECK(cudaFree(d_y));

  return 0;
}

性能表现

nvcc  add.cu -o add_cu -gencode arch=compute_75,code=sm_75
time ./add_cu

GPU Kernel 'add' execution time: 48.6738 ms
Max error: 0

real 0m19.413s
user 0m15.308s
sys 0m4.014s
  • 计算耗时: GPUKernel函数的执行耗时仅为 48.7毫秒。
  • 总耗时: 程序总耗时为 19.4秒。

性能分析

单看核心计算任务，GPU (48.7ms) 的速度是CPU (3740ms) 的 约75倍。这完美体现了GPU在处理数据并行任务时的绝对优势。CPU需要串行执行10亿次加法(此处只考虑单核场景)，而GPU则将任务分配给成千上万个线程同时处理。

但是虽然GPU计算本身极快，但程序的总耗时 (19.4s) 却和CPU版本 (21.4s) 相差无几。这是为什么呢？主要是CPU和GPU通讯的开销。这里下一篇文章会详细介绍。

编译-Fat Binary

nvcc add.cu -o add_cu -gencode arch=compute_75,code=sm_75 上面的例子中，我们看到这个编译指令。add.cu被编译为二进制文件add_cu。它具体是怎么做的呢？

  1. 主机代码编译: 将C/C++代码（在CPU上运行的部分）交由系统的主机编译器（如GCC、MSVC）编译成标准的CPU目标代码。
  2. 设备代码编译: 将在__global__函数（如add）中定义的GPU代码，编译成两种主要格式：
  • SASS (Streaming Assembler): 这是特定GPU架构的原生机器码。例如，为NVIDIA T4 GPU (Turing 架构 代号 sm_75 ) 编译的SASS，只能在该架构上最高效地运行。NVCC可以为多种指定的架构预编译多份SASS代码。
  • PTX (Parallel Thread eXecution): arch=compute_75 指示编译器生成一份 PTX 代码，确保程序能在任何不低于 Turing 架构的新 GPU 上通过 JIT 编译运行（向前兼容性）。

这两种设备代码连同主机代码一起，被打包进一个可执行文件中，形成所谓的胖二进制 (Fat Binary)。它“胖”在包含了一份主机代码和多份针对不同GPU架构的设备代码。

程序加载 - cubin loading

程序启动

操作系统加载可执行文件，CPU 开始执行主机代码。

首次 CUDA 调用

当代码第一次调用任何 CUDA API 函数时（比如 cudaSetDevice, cudaMalloc，或者第一个Kernel函数启动），CUDA 运行时库 (CUDA Runtime Library) 会被初始化。

此处就是所谓的GPU上下文初始化/CUDA上下文初始化，主要步骤：

1. 硬件准备与唤醒

  • 从低功耗的待机模式唤醒，进入高性能的计算模式；
  • 加载驱动模块（如NVIDIA CUDA Driver或AMD ROCm），并检测可用GPU设备及其属性（如显存大小、计算能力、NVLink连接）。

2. CUDA上下文数据结构创建

CPU侧创建上下文信息的数据结构：创建一个统一虚拟地址空间（UVA)，这个空间可以将所有的系统内存和所有GPU的内存都映射进来，共享一个单一的虚拟地址空间。(每次cudaMalloc都会增加一条记录)

3. 特定GPU上创建上下文

  • 在显存中为当前进程分配并建立页表结构
    a. NVIDIA驱动程序（在CPU上）查询其内部维护的、用于管理GPU物理显存的数据结构（即VRAM Allocator，跨进程维护），以找到一个空闲的物理地址。CPU本地软件操作，不涉及与GPU的硬件通信。
    b. CPU在自己的内存（RAM）里，准备好了要写入的数据内容；
    c. NVIDIA驱动程序（在CPU上）命令DMA引擎将对应数据复制到显存；
  • 分配Pinned Memory命令缓冲区
  • 通过MMIO配置GPU的MMU硬件(PMMU 控制寄存器)，告诉它页表的起始位置

4. 上下文就绪

上下文完全建立，后续的Kernel函数启动、内存拷贝等命令可以通过流 (Stream) 机制提交到其命令缓冲区，由GPU异步执行。

首次调用add<<<...>>>()时，进行Kernel函数进行加载

1. 检测硬件

它会查询当前的 GPU，识别出具体架构。

2. 寻找最佳匹配 (SASS)

然后，它会在 Fat Binary 的设备代码段中进行搜索，寻找有没有预编译好的、针对 sm_75 的 SASS 代码。

3. 没有找到完全匹配的 SASS 代码

如果没有找到完全匹配的 SASS 代码运行时会找到 PTX 中间代码，并调用集成在 GPU 驱动中的 JIT (Just-In-Time) 编译器将其即时编译(JIT)为目标GPU的SASS代码; (cpu上完成)；

为了避免每次运行程序都重新进行 JIT 编译，NVIDIA 驱动通常会缓存 JIT 编译的结果。NVIDIA驱动会在用户的home目录下创建一个计算缓存，通常是 ~/.nv/ComputeCache。

4. cubin loading (cubin 是 CUDA binary 的缩写)

a. 将准备好的 SASS 代码（无论是来自 Fat Binary 还是 JIT 编译的结果）申请显存空间；通过DMA复制到显存；

b. 驱动程序在其内部的表格中，将Kernel函数 add 与其在 VRAM 中的地址关联起来。后续调用 add<<<...>>>() 时，运行时会将一个包含该 VRAM 地址的启动命令提交到流中，由 GPU 异步执行

程序执行 - Kernel Launch

一个常见的误解是CPU会直接、实时地控制GPU。实际上，考虑到CPU和GPU是两个独立的处理器，并且通过PCIe总线连接，直接的、同步的控制会带来巨大的延迟和性能开销。因此，现代GPU采用了一种高效的异步通信模型，其核心就是 命令缓冲区（Command Buffer）与门铃（Doorbell）机制。这也是CUDA Streaming的底层通讯机制。

Command Buffer + Doorbell 机制

  1. cpu先把需要执行的命令写到ring buffer命令缓冲区(Pinned Memory,位于主机内存); 更新w_ptr
  2. 在适当的时候通过MMIO设置Doorbell Register，告诉GPU有新任务需要处理
  3. GPU上的DMA引擎将ring buffer命令缓冲区[r_ptr, w_ptr)复制到显存中，然后开始执行；(其中w_ptr和r_ptr可以理解为相对于 Ring Buffer 基地址 (Base Address) 的偏移量)

下面对于部分由代表型的API的执行逻辑进行单独阐述。

CPU 执行到cudaMalloc

cudaMalloc 是一个同步阻塞调用，它不使用上述的流式命令缓冲区机制。（CUDA 11.2+支持cudaMallocAsync可实现异步分配）

  1. CPU 线程调用 cudaMalloc()。CUDA 运行时库将此请求转发给 NVIDIA 驱动程序
  2. 驱动程序向物理VRAM Allocator请求物理内存，向 UVA Manager 请求虚拟地址，更新UVA映射表；（物理VRAM Allocator是跨进程的，维护整个GPU 物理显存的使用情况）
  3. 更新 GPU page table[Command Buffer + Doorbell方式,特定的、高优先级的通道，非默认的Stream]，刷新TLB
  4. 返回虚拟内存指针

与malloc的不同之处

  1. Lazy Allocation vs. Eager Allocation
    malloc支持overcommit，实际物理内存的分配发生在访问时(Lazy Allocation)，通过缺页中断（Page Fault）按需映射到物理内存；而cudaMalloc是同步分配连续的物理显存(Eager Allocation)，保证了后续使用的确定性，但初始开销更高。
  2. system call overhead
    cudaMalloc直接陷入内核，调用GPU驱动分配物理内存；而malloc本身是C库函数(用户态), 向操作系统“批发”大块内存，然后在用户程序请求时“零售”出去。避免内存分配时昂贵的系统调用和缺页异常开销
  • 申请<128KB内存时，会优先在freelist中查找是否有合适的空闲 Chunk，没有找到，才会通过brk系统调用向操作系统申请内存
  • 申请>=128KB内存时，会直接通过mmap系统调用向操作系统申请内存，free时也会直接释放
  1. 释放策略
    cudaFree会直接释放，而free对于brk/sbrk分配的内存不会直接释放(物理内存和虚拟内存都不释放，为了避免Page Fault引入的性能开销就没有释放物理内存)，用户态维护freelist，同时会合并连续空闲的虚拟地址空间，有效减少内存碎片(coalescing)。

CPU 执行到 cudaMemcpy、cudaMemset

通过Command Buffer + Doorbell 机制提交命令到GPU; 然后同步或者异步等待

CPU 执行到Kernel函数add<<<...>>>()

CPU侧：命令打包与提交

  1. 驱动将Kernel函数启动所需信息打包成一个命令
    命令包括：启动Kernel函数，Kernel函数对应的add SASS 代码的入口地址，执行配置(Grid 维度、Block 维度、共享内存大小等)、参数指针(GPU虚拟地址)
  2. 将命令写入主机端的 Pinned Memory Ring Buffer
  3. 通过 MMIO 写 Doorbell 寄存器，通知 GPU

GPU侧: 命令获取与运行

  1. 通过 DMA 从 Pinned Memory 读取Ring buffer部分内容
  2. 命令解码
  • GPU 的命令处理器 (Front-End) 从其内部队列中取出命令包。
  • 它开始解码这个命令包，识别出这是一个“Kernel函数启动”任务，并解析出所有的执行参数（Grid/Block 维度、Kernel函数地址等）。
  1. 工作分发
    int blockSize = 256;
    int numBlocks = (N + blockSize - 1) / blockSize;
    add<<<numBlocks, blockSize>>>(N, d_x, d_y);

  • 命令处理器根据 Grid 的维度，将整个计算任务分发成一个个独立的Thread Blocks。
  • GPU的全局调度器（GigaThread Engine）,将Thread Blocks分配给有空闲资源的 SM。一个线程块从生到死都只会在一个 SM 上执行，不会迁移。
  1. 线程块调度与执行
    __global__
    void add(int n, float *x, float *y)
    {
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    if (index < n) {
    y[index] = x[index] + y[index];
    }
    }

  • 每个 SM 接收到一个或多个线程块，SM 内部的硬件调度器 (Scheduler)进一步将每个线程块内部的线程，按照threadIdx的顺序，每 32 个线程划分成一个 Warp
  • 。比如，一个有 256 个线程的线程块，会被划分为 8 个 Warps (Warp 0: 线程 0-31, Warp 1: 线程 32-63, ...)。
  • SM 内部的硬件调度器 (Scheduler) Warps分配给 SM 内的CUDA Cores 和其他执行单元（如 Tensor Cores）去执行。
  • CUDA 核心开始执行位于指定 SASS 地址的机器指令，进行实际的计算。
  1. 完成与资源回收
  • 当一个线程块完成了所有计算，它所占用的 SM 资源（如寄存器、共享内存）会被释放，SM 可以接收新的线程块。
  • 当整个 Grid 的所有线程块都执行完毕，这个Kernel函数启动任务就算完成了。

Grid、Thread Block、Warp、Thread、SM这些概念到底是干啥的。下面结合GPU的硬件架构详细介绍。

GPU的硬件架构

如上是NVIDIA GA100 GPU的架构图：

A100 GPU 架构图

计算单元

GPC

Graphics Processing Cluster, 一个GPU包含多个GPC, 一个GPC包含多个TPC

TPC

Texture Processing Cluster, 一个TPC包含多个SM

SM

Streaming Multiprocessor, SM是GPU执行计算任务的核心单元，它是

  • CUDA Cores (执行FP32/INT32等通用计算的ALUs/FPUs)
  • Tensor Cores
    一个硬件单元，专门处理**FMA(Fused Multiply-Add)**操作，能在一个时钟周期内完成一个小的矩阵乘加运算（一个4x4的FP16矩阵相乘后累加到另一个4x4矩阵上）
    深度学习绝大部分的计算都是FMA操作，NVidia工程师为此专门设计专用计算单元。
  • 寄存器 (Register File)、共享内存 (Shared Memory)
  • L1数据缓存/指令缓存 (L1 Data Cache / Instruction Cache)
  • Warp调度器 (Warp Scheduler) 等关键组件

单个SM的架构图如下：

接口

  • PCIe 负责CPU与GPU的通讯，DMA模式
  • NVLINK 负责GPU间的通讯

内存与缓存

其中HBM和L2 Cache是整个GPU共享的;

而L1 Cache/Shared Memory则是SM维度独享的;

Shared Memory是每个SM内部的一块高速、可编程的片上缓存。同一线程块（Block）内的所有线程都可以访问它，速度远快于访问全局显存（HBM）。它是实现Block内线程高效协作和数据交换的核心，对于矩阵乘法等需要数据复用的算法至关重要。

速度由快到慢依次为 寄存器 -> L1 Cache -> L2 Cache -> HBM -> DRAM(主机内存)

编程模型 vs 硬件执行模型

编程模型

将一个待批量并发的数据组织成Grid、Thread Block、Thread的结构。

Grid和Thread Block可以是1维的也可以是2维或者3维的。这里这么设计，感觉主要是为了让程序员可以根据实际处理的结构能够更自然的思考,同时可以覆盖数据局部性需求，比如，我要处理一个1维数据，自然的我们就可以把Grid和Thread Block定义为1维的。比如上面例子中计算1维数组的加法，就可以用1维的Grid和Thread Block。

int blockSize = 256;
int numBlocks = (N + blockSize - 1) / blockSize;
add<<<numBlocks, blockSize>>>(N, d_x, d_y);

__global__
void add(int n, float *x, float *y)
{
  int index = blockIdx.x * blockDim.x + threadIdx.x;
  if (index < n) {
    y[index] = x[index] + y[index];
  }
}

Grid视图：

这行代码是CUDA编程的基石(SIMT)，它将软件层面的线程坐标映射到数据上的全局索引。

  • threadIdx.x: 当前 Thread 在其 Block 内的 x 坐标。范围是 0 到 blockDim.x - 1。
  • blockDim.x: 每个 Block 在 x 维度上有多少个 Thread。（在我们例子中是256）。
  • blockIdx.x: 当前 Block 在 Grid 中的 x 坐标。范围是 0 到 gridDim.x - 1。
  • gridDim.x: Grid 在 x 维度上有多少个 Block。

blockIdx.x * blockDim.x计算出了当前线程块之前所有线程块包含的线程总数（偏移量），再加上threadIdx.x，就得到了当前线程在整个Grid中的全局唯一ID。这保证了10亿个元素，每个都能被一个特定的线程处理到。

这里解释下上面提到的数据局部性: y[index] = x[index] + y[index]; 可以合并访存 (Coalesced Memory Access)。即一个Warp中的32个线程访问连续的32个内存地址，GPU硬件可以将其合并成一次或少数几次宽内存事务，极大提升访存效率。

而当我们要处理一个二维矩阵或图像时，最自然的思考方式就是二维的。这时候我们可以用2维的Grid和Thread Block。

dim3 blockSize(16, 16); // 16x16 = 256 线程/块
dim3 gridSize((N + blockSize.x - 1) / blockSize.x, (N + blockSize.y - 1) / blockSize.y);

__global__ void matrixMulGPU(const float* A, const float* B, float* C, int N) {

    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    if (row < N && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < N; ++k) {
            sum += A[row * N + k] * B[k * N + col];
        }
        C[row * N + col] = sum;
    }
}

Grid视图：

硬件层面

将整个GPU的运算单元分为 GPU、SM、Warp和Core。

软件层面将grid切分成多个Thread Block是为了对硬件的抽象，这样程序员就不必关心GPU具体有多少个物理核心、多少个SM。

Thread Block是最小的“资源分配与调度”单位，Warp是最小的硬件调度单位。

所以整个编程模型大概就是：

一个任务软件层面上被分为Grid和Thread Block，Thread Block被分配给硬件的SM，SM又将Thread Block按照32个Thread为一组分成Warp，分配给Warp scheduler执行。

最终的视图大概是这样的：

隐藏延迟 - hide latency

前面已经看到一个计算任务对应一个Grid，一个Grid又由多个Thread Block组成，GPU的全局调度器（GigaThread Engine）将Thread Blocks分配给有空闲资源的 SM。(多个Thread Blocks可以被分配给一个SM,取决于共享内存、寄存器使用的使用情况)

一个Thread Block被分解成多个Warp（例如，一个1024线程的Block被分解成32个Warp)。SM内部的调度硬件，会将这32个Warp分配给它内部的4个Warp Scheduler。通常会尽量均匀分配，比如每个Warp Scheduler分到8个Warp。

而一个Warp Scheduler同一时刻只能运行一个Warp, 当某个正在执行的Warp因为等待内存而暂停时，它可以立刻从剩下的Warp中挑选一个就绪的来执行。这就是所谓的隐藏延迟 (hide latency)。而如何充分利用这个特性呢？给每个Warp Scheduler足够多的可切换的Warp。

每个SM都包含一个巨大、单一的物理寄存器文件，为实现零开销Warp上下文切换的提供了硬件基础。这是与CPU昂贵的上下文切换（需要保存和恢复大量状态）的根本区别。

要让每个 Warp Scheduler (Warp 调度器) 有足够的可切换 Warp，其本质是提高 GPU 的占用率。占用率指的是一个 SM 上实际活跃的 Warp 数量与该 SM 理论上能支持的最大 Warp 数量的比例。

一个 SM 能同时运行多少 Warp（**一个 SM 在同一时刻只能为一个 Kernel 服务，但可以同时运行该Kernel的多个线程块(只要资源允许)**），取决于以下三个主要资源的限制：

1. Registers

每个线程都需要使用寄存器来存储其局部变量。一个 SM 上的寄存器总数是固定的

假设一个 SM 有 65536 个寄存器，最大支持 2048 个线程 (64 Warps)。 每个Kernel需要 64 个寄存器，那么一个 Block (假设 256 线程) 就需要 256 * 64 = 16384 个寄存器。这个 SM 最多可以容纳 65536 / 16384 = 4 个这样的 Block，也就是 1024 个线程 (32 Warps)，占用率为 50%。如果 Kernel 每个线程需要 128 个寄存器，那么这个 SM 只能容纳 2 个这样的 Block，占用率就更低了。

2. Shared Memory

共享内存是分配给每个线程块 (Block) 的、速度很快的片上内存。一个 SM 上的共享内存总量是固定的。

假设一个 SM 有 96KB 共享内存，最大支持 16 个 Block。如果Kernel 每个 Block 需要 32KB 共享内存，那么这个 SM 最多只能同时运行 96KB / 32KB = 3 个 Block。在这个场景下，共享内存成为了主要的限制因素。这就将 SM 上并发的 Block 数量上限从硬件支持的 16 个锐减到了 3 个，从而严重限制了 SM 上的总并发 Warp 数量，降低了占用率。

3. 线程块/线程数限制

每个 SM 架构本身就有硬件限制，比如一个 SM 最多能同时调度多少个 Block（例如 16 或 32），以及最多能同时管理多少个线程（例如 2048）。这个是硬性上限，无法通过代码改变。

不过提高 GPU 的占用率来隐藏延迟也不是万能的，隐藏延迟的有效性，本质上取决于 Warp调度器是否有“就绪态”的Warp可供切换。比如：如果一个Kernel非常简单，每个线程只使用极少的寄存器，并且不使用共享内存，那么一个SM上可能会驻留大量的Warp。但如果这个Kernel的计算是访存密集型且延迟很高的，同时计算/访存指令比例很低，那么即使占用率达到100%，Warp调度器可能依然会“无Warp可调”，因为所有Warp都在等待数据返回。这时候我们就不得不提另外一个概念，访存比(Ratio = Total Bytes / Total FLOPs)或者计算强度(Roofline,I = Total FLOPs / Total Bytes), 说白了，就是看一个程序是计算密集型(Compute-bound)还是IO(内存访问)密集型(Memory-bound)。可以使用NVIDIA Nsight Compute分析Kernel函数的占用率和计算强度。 不过这里不做延伸了，放到下篇性能优化中讲。

SIMD vs SIMT

前面CUDA Demo中我们已经知道Kernel函数add会被启动成茫茫多的线程执行,每个线程通过计算 blockIdx 和 threadIdx 来处理不同的数据。

__global__
void add(int n, float *x, float *y)
{
  int index = blockIdx.x * blockDim.x + threadIdx.x;
  if (index < n) {
    y[index] = x[index] + y[index];
  }
}

从程序员的角度看，我们似乎是在编写多线程（Multiple Threads）程序。但从硬件的角度看，它是如何让这么多线程同时执行同一条指令（Single Instruction）的呢？

这种 "单指令，多线程"（Single Instruction, Multiple Threads, SIMT）的编程模型，正是CUDA的魅力所在。SIMT通过线程编程模型巧妙的隐藏了底层SIMD的执行细节。而要理解SIMT，就不得不提在CPU中广泛使用的SIMD技术。

SIMD(Single instruction, multiple data)

在传统的标量计算模型中，CPU的一条指令一次只能操作单个数据。例如，一次浮点加法就是double + double;

当处理如图形、音频或科学计算中常见的大规模数据集时，这种“一次一个”的模式效率极低，因为我们需要对海量数据重复执行完全相同的操作，这暴露了标量处理的瓶颈。

展开阅读全文

本文系作者在时代Java发表，未经许可，不得转载。

如有侵权，请联系nowjava@qq.com删除。

编辑于
举报

关注时代Java

教程文章

文章订阅

关注时代Java

Copyright © 2014-2024NowJava. All Rights Reserved. 时代Java 版权所有 京ICP备14059627号-7

关注时代Java