文章

gpu-glossary

Posted 2026-03-19 Updated 2026-03- 19
By herbduan
167~215 min read

源网站：

https://modal.com/gpu-glossary/readme

机器翻译于 2026-03-12

我们在 Modal 使用 GPU 的过程中，发现相关文档资料非常零散，导致很难将技术栈中不同层级的概念（例如流式多处理器架 (Streaming Multiprocessor Architecture)、计算能力 (Compute Capability) 和 nvcc 编译器标志）关联起来。为了解决这个问题，我们撰写了这个术语表。

为此，我们阅读了 NVIDIA 的 PDF 文档，潜心研究了相关的 Discord 优质社群，甚至购买了实体教科书，最终汇集了这份涵盖整个技术栈的术语表。

与 PDF、Discord 或书籍不同，这份术语表是一份超文本文档——所有页面都相互链接。因此，你可以随时跳转去了解线程束调度器 (Warp Scheduler)，从而更好地理解你在CUDA C++ 编程模型文章中遇到的线程 (threads)。

你也可以按顺序阅读。请使用方向键、每页底部的箭头或目录（在桌面端的侧边栏或移动端的汉堡菜单中）在页面之间导航。

本术语表的源代码可在 GitHub 上获取。

正在基于 GPU 进行构建？我们对此颇有心得。

Modal 是一个围绕全球 GPU 集群构建的、符合人体工程学的 Python SDK。你无需担心配额申请、驱动程序兼容性问题或管理庞大的机器学习依赖，即可即时部署无服务器 AI 工作负载。

在 GPU 上部署

设备硬件 (Device Hardware)

GPU 硬件是运行计算工作负载的物理组件。

CUDA（设备架构）

CUDA 设备架构是指 GPU 的物理设计，特别是其计算核心的组织方式。

历史背景在引入 CUDA 之前，GPU 主要用于图形处理。它们包含专门的硬件单元，用于执行渲染管线的特定阶段，例如顶点着色、几何着色和像素着色。这种固定功能管线（fixed-function pipeline）虽然高效，但缺乏灵活性。

在 G71 架构（在 GeForce 7 系列 GPU 中使用）中，着色器（绿色）是独立的、特定于任务的单元。

2006 年，NVIDIA 推出了 Tesla 架构，这是他们的第一个统一计算架构。它用一个由多个相同核心组成的阵列取代了固定功能管线，这些核心能够执行通用的计算任务。这种设计（在 G80 芯片中首次亮相）不仅更适用于图形处理，还为通用 GPU（GPGPU）计算铺平了道路，并催生了 CUDA。

G80 芯片是首款采用 NVIDIA Tesla 统一计算架构的芯片。它拥有 128 个流处理器（stream processors），这些是现代 CUDA 核心的前身。

流式多处理器 (Streaming Multiprocessor, SM)

当我们为 GPU 编程时，我们实际上是在为流式多处理器（Streaming Multiprocessors）生成指令序列来执行。

H100 GPU 的流式多处理器内部架构图。GPU 核心显示为绿色，其他计算单元为栗色，调度单元为橙色，内存为蓝色。修改自 NVIDIA 的 H100 白皮书。

NVIDIA GPU 的流式多处理器（SM）大致相当于 CPU 的核心。也就是说，SM 既执行计算，也在寄存器中存储可用于计算的状态，并配有关联的缓存。与 CPU 核心相比，GPU 的 SM 是简单且能力较弱的处理器。SM 内的执行是指令流水线式的（instruction pipelined）（几乎所有自 20 世纪 90 年代以来的 CPU 都如此），但没有推测执行（speculative execution）或指令指针预测（instruction pointer prediction）（与所有当代高性能 CPU 不同）。

然而，GPU SM 可以并行执行更多的线程 (threads)。

作为对比：一个 AMD EPYC 9965 CPU 的功耗最高为 500W，拥有 192 个核心，每个核心最多能同时为两个线程执行指令，总共可并行 384 个线程，每个线程约消耗 1.25W。

一个 H100 SXM GPU 的功耗最高为 700W，拥有 132 个 SM，每个 SM 有四个线程束调度器 (Warp Schedulers)，每个调度器可以在每个时钟周期内向 32 个线程（即一个线程束 (warp)）并行发出指令，总计 128 × 132 > 16,000 个并行线程，每个线程约消耗 5cW。请注意，这是真正的并行：每 16,000 个线程中的每一个都能在每个时钟周期取得进展。

GPU SM 还支持大量的并发线程——这些线程的指令是交错执行的。

H100 上的单个 SM 可以并发执行多达 2048 个线程，这些线程被分成 64 个线程组，每组 32 个线程。对于 132 个 SM，总共可以并发执行超过 250,000 个线程。

CPU 也可以并发运行许多线程。但是，线程束之间的切换速度仅为单个时钟周期（比 CPU 上的上下文切换快 1000 倍以上），这同样得益于 SM 的线程束调度器。可用线程束的数量和线程束切换的速度有助于隐藏延迟 (hide latency)，这些延迟由内存读取、线程同步或其他昂贵指令引起，从而确保 CUDA 核心 (CUDA Cores) 和张量核心 (Tensor Cores) 提供的算术带宽 (arithmetic bandwidth) 得到充分利用。

这种延迟隐藏 (latency-hiding) 是 GPU 强大的秘诀。CPU 试图通过维护大型、由硬件管理的缓存和复杂的指令预测来向终端用户和程序员隐藏延迟。这些额外的硬件限制了 CPU 可分配给计算的芯片面积、功耗和散热预算的比例。

与 CPU 相比，GPU 将更多的芯片面积用于计算（绿色），而较少用于控制和缓存（橙色和蓝色）。修改自 Fabien Sanglard 的博客中的图表，该图表可能又修改自 CUDA C 编程指南。

对于像神经网络推理或顺序数据库扫描这样的程序或函数，程序员可以相对直接地表达 缓存的行为——例如，存储每个输入矩阵的一块数据并将其保存在缓存中足够长的时间以计算相关的输出——其结果是更高的吞吐量。

核心 (Core)

GPU 核心是 GPU 中执行计算的最小单位。核心通常针对特定类型的计算任务进行优化。

H100 GPU 的流式多处理器内部架构。CUDA 核心和张量核心显示为绿色。修改自 NVIDIA 的 H100 白皮书。

GPU 核心类型的例子包括：

  • CUDA 核心 (CUDA Core)：执行标准的单精度浮点（FP32）或整数（INT32）算术运算。

  • 张量核心 (Tensor Core)：专门用于执行矩阵乘加（MMA）运算，这是深度学习工作负载中的常见操作。

特殊功能单元 (Special Function Unit, SFU)

特殊功能单元（SFU）是一种 GPU 核心，专门用于执行 CPU 上不常见的超越函数（transcendental functions）。

H100 SM 的内部架构。特殊功能单元以粉红色显示，与加载/存储单元一起。修改自 NVIDIA 的 H100 白皮书。

例如，sin(x)、cos(x)、log2(x) 和 1/x 等函数的计算通常由 SFU 处理。这些单元在图形和科学计算应用中尤为重要，因为这些应用经常依赖这些数学函数。

加载/存储单元 (Load/Store Unit, LSU)

加载/存储单元（LSU）是一种 GPU 核心，专门用于处理内存操作。

H100 SM 的内部架构。加载/存储单元与特殊功能单元一起显示为粉红色。修改自 NVIDIA 的 H100 白皮书。

最重要的是，对于面向吞吐量的计算，LSU 负责将数据从 GPU 的全局内存移动到 SM 的寄存器文件中，供算术核心使用。

线程束调度器 (Warp Scheduler)

线程束调度器（Warp Scheduler）是流式多处理器（SM）内的一个组件，负责管理线程束 (warps)的执行。

H100 SM 的内部架构。线程束调度器和分发单元显示为橙色。修改自 NVIDIA 的 H100 白皮书。

这些线程 (threads) 的分组，即线程束 (warps)，由线程束调度器进行调度。每个时钟周期，线程束调度器都会选择一个准备好执行的线程束，并向该线程束中的线程发出下一条指令。

通过在等待内存操作或其他长延迟事件的线程束之间快速切换，线程束调度器在隐藏延迟 (latency hiding) 中扮演着关键角色。这种快速的上下文切换能力确保了 SM 的执行单元能够持续保持繁忙，从而最大化吞吐量 (throughput)。

CUDA 核心 (CUDA Core)

CUDA 核心是 NVIDIA GPU 中的基本计算单元。它们是 GPU 的主力，负责执行构成任何应用程序的大部分算术和逻辑运算。

H100 SM 的内部架构。CUDA 核心和张量核心以绿色显示。请注意张量核心的尺寸更大但数量更少。修改自 NVIDIA 的 H100 白皮书。

它们与张量核心 (Tensor Cores) 形成对比，后者是专门用于特定矩阵运算的计算单元。

从历史上看，CUDA 核心被称为“流处理器（streaming processors）”，这反映了它们在处理数据流方面的作用。虽然“核心”这个词让人联想到 CPU 核心，但重要的是要认识到 CUDA 核心要简单得多，数量也更多。一个高端 CPU 可能有几十个核心，而一个现代 GPU 可能有数千个 CUDA 核心。

张量核心 (Tensor Core)

张量核心是专门的 GPU 核心，旨在加速矩阵乘加（matrix multiply-accumulate, MMA）运算，这是深度学习和科学计算中常见的操作。

H100 SM 的内部架构。请注意张量核心的尺寸更大但数量更少。修改自 NVIDIA 的 H100 白皮书。

在单条指令中操作更多数据，可以显著降低功耗并提高算术带宽 (arithmetic bandwidth)。

张量核心在单个时钟周期内对 4x4 矩阵执行 D = A @ B + C 形式的运算。输入矩阵 A 和 B 通常是 FP16 或 BF16 格式，而累加矩阵 C 和 D 可以是 FP16、BF16 或 FP32 格式。

一些现代张量核心，例如在 H100 GPU 中找到的那些，还支持 FP8 和 INT8 格式。

张量核心与线程束 (warp) 内的线程 (threads) 协同工作，以高效处理更大的矩阵。通过将较大的矩阵分块并将其分配给线程束内的各个线程，GPU 可以并行执行多个 4x4 MMA 运算。

张量核心 MMA 中用于 C = A @ B 的寄存器使用情况。R11、R17、R16 和 R18 寄存器分别用于指令 1、2、3 和 4。详见周围文本。

张量内存加速器 (Tensor Memory Accelerator, TMA)

张量内存加速器（TMA）是一个专用硬件单元，旨在加速张量核心 (Tensor Cores) 的内存操作。

H100 流式多处理器（SM） 的内部架构。请注意，张量内存加速器位于 SM 的底部，由四个子单元共享。修改自 NVIDIA 的 H100 白皮书。

TMA 从全局内存加载数据并将其存储在共享内存中，使其可以被 SM 内的张量核心高效访问。这种异步数据传输允许张量核心在数据加载的同时执行计算，从而减少了内存访问的延迟 (latency) 并提高了整体吞吐量 (throughput)。

流式多处理器架构 (Streaming Multiprocessor Architecture)

流式多处理器架构（SM 架构）是指 GPU 的流式多处理器（SM）的特定设计和组织。

Tesla 架构 SM 的框图。它包含 8 个流处理器（SP）（现代 CUDA 核心的前身）、2 个特殊功能单元（SFU）、指令和数据缓存以及共享内存。修改自 NVIDIA 的 Tesla 技术简介。

不同的 GPU 家族具有不同的 SM 架构，例如：

  • Tesla (2006)

  • Fermi (2010)

  • Kepler (2012)

  • Maxwell (2014)

  • Pascal (2016)

  • Volta (2017)

  • Turing (2018)

  • Ampere (2020)

  • Hopper (2022)

  • Ada Lovelace (2022)

  • Blackwell (2024)

每个新的 SM 架构通常都会带来改进和新功能，例如更多的核心、更大的缓存、新的指令类型和更高的能效。这些进步与 GPU 的计算能力密切相关，后者定义了 GPU 支持的特性和能力。

纹理处理集群 (Texture Processing Cluster, TPC)

纹理处理集群（TPC）是 GPU 内流式多处理器（SM）的一种组织单元。

在现代 NVIDIA GPU 中，一个 TPC 通常包含一个或多个 SM，以及专门用于纹理处理的硬件。TPC 的具体配置（例如每个 TPC 的 SM 数量）因不同的 GPU 架构而异。例如，在 Ada Lovelace 架构中，每个 TPC 包含两个 SM。

图形/GPU 处理集群 (Graphics/GPU Processing Cluster, GPC)

图形处理集群（GPC）是 GPU 内纹理处理集群（TPC）的一种组织单元。

一个 GPC 通常包含多个 TPC。与 TPC 一样，GPC 的具体配置因不同的 GPU 架构而异。

例如，一个完整的 AD102 核心（用于 RTX 4090）有 12 个 GPC，每个 GPC 有 6 个 TPC，每个 TPC 有 2 个 SM，总计 144 个 SM。

寄存器文件 (Register File)

寄存器文件是位于每个流式多处理器（SM）上的一个小而高速的内存。它为在 SM 上执行的线程 (threads) 提供寄存器。

H100 SM 的内部架构。寄存器文件显示为蓝色。修改自 NVIDIA 的 H100 白皮书。

寄存器文件被划分为多个部分，供 SM 上并发执行的线程束 (warps) 使用。每个线程都有自己的一组专用寄存器，用于存储其私有数据。

访问寄存器文件非常快，通常只需要一个时钟周期。这使得寄存器成为存储频繁访问数据（如循环计数器和临时变量）的理想选择。

L1 数据缓存 (L1 Data Cache)

L1 数据缓存是位于每个流式多处理器（SM）上的一个小而高速的内存。它用于存储频繁访问的数据，减少从较慢的内存层级（如全局内存）中获取数据的需求。

H100 SM 的内部架构。L1 数据缓存以浅蓝色显示。修改自 NVIDIA 的 H100 白皮书。

每个 SM 将该内存划分为 L1 数据缓存和共享内存。这种划分可以由程序员在每次核函数 (kernel) 启动时进行配置。

张量内存 (Tensor Memory)

我们找不到这个术语的任何好的参考资料！

如果你知道它是什么，请在 GitHub 上告诉我们！

GPU RAM

GPU RAM，也称为图形 RAM 或视频 RAM (VRAM)，是 GPU 用于存储其计算所需数据的高速内存。

高带宽内存（HBM）是一种常用于现代 GPU 的堆叠式 RAM 技术。它提供了比传统 DDR SDRAM 更高的带宽和更低的功耗。

与 CPU 使用的系统 RAM 不同，GPU RAM 经过优化，可实现高带宽和并行访问，这对于处理大规模数据集和执行并行计算至关重要。

GPU RAM 对于 GPU 性能至关重要，因为它直接影响 GPU 访问和处理数据的速度。GPU RAM 的容量越大、速度越快，GPU 就能越快地执行复杂计算，从而在游戏、科学模拟和机器学习等应用中获得更好的性能。

GPU RAM 存储着CUDA 编程模型的全局内存。

设备软件 (Device Software)

GPU 设备软件是指在 GPU 硬件上直接执行的程序和指令。

CUDA（编程模型）

CUDA 编程模型是一种用于在 NVIDIA GPU 上进行并行计算的软件抽象。它允许开发者将计算任务分解为可以在 GPU 的众多核心上并行执行的较小子任务。

CUDA 编程模型的线程层级结构，从单个线程到线程块，再到线程块网格（左侧），映射到硬件上，从 CUDA 核心到流式多处理器，再到整个 GPU（右侧）。修改自 NVIDIA 的 CUDA Refresher: The CUDA Programming Model 和 NVIDIA CUDA C++ 编程指南中的图表。

CUDA 编程模型的核心概念包括：

  • 核函数 (Kernel)：一个在 GPU 上执行的函数，由大量线程并行运行。

  • 线程 (Thread)：执行计算的最小单位。

  • 线程块 (Thread Block)：一组可以相互协作的线程，例如通过共享内存和同步。

  • 网格 (Grid)：一组线程块，它们共同执行一个核函数。

该模型还定义了一个内存层级，包括每个线程的私有寄存器、每个线程块的共享内存，以及所有线程均可访问的全局内存。

流式汇编器 (Streaming ASSembler, SASS)

SASS 是 NVIDIA GPU 的原生指令集。它是一种低级汇编语言，直接由 GPU 的流式多处理器（SM）执行。

当使用 CUDA C++ 或其他高级语言编写 CUDA 核函数 (kernel) 时，nvcc 编译器首先将代码编译成 PTX，这是一种中间表示。然后，在运行时，GPU 驱动程序将 PTX 代码即时（JIT）编译成目标 GPU 架构的 SASS 代码。

SASS 指令针对特定的 GPU 架构进行了高度优化，可以直接控制 GPU 的硬件资源。这使得 SASS 成为实现最高性能的关键，但也使其编写和维护变得非常具有挑战性。

并行线程执行 (Parallel Thread eXecution, PTX)

PTX 是 NVIDIA GPU 的一种中间表示（IR）。它是一种虚拟机和指令集架构（ISA），为 GPU 提供了一个独立于硬件的编程模型。

PTX 的机器模型。修改自 PTX ISA 参考手册。

当使用 CUDA C++ 编写 CUDA 核函数 (kernel) 时，nvcc 编译器首先将代码编译成 PTX。然后，GPU 驱动程序将 PTX 代码即时（JIT）编译成目标 GPU 架构的原生 SASS 代码。

这种两步编译过程允许 CUDA 程序在前向兼容新的 GPU 架构。只要新 GPU 的驱动程序支持，为旧 GPU 编译的 PTX 代码就可以在新 GPU 上运行。

计算能力 (Compute Capability)

计算能力是 NVIDIA 用来描述其 GPU 硬件特性的版本号。它由一个主版本号和一个次版本号组成，例如 7.5、8.0 或 9.0。

主版本号表示 GPU 的核心流式多处理器架构。例如，计算能力 7.x 对应于 Volta 架构，8.x 对应于 Ampere 架构，9.x 对应于 Hopper 架构。

次版本号表示对核心架构的增量改进，例如更大的缓存、新的指令或更高的时钟速度。

计算能力决定了 GPU 支持的功能，例如：

线程 (Thread)

在 CUDA 编程模型中，线程是执行计算的最小单位。

CUDA 编程模型中的线程层级结构。修改自 NVIDIA 的 CUDA Refresher 和 CUDA C++ 编程指南。

当一个 CUDA 核函数 (kernel) 启动时，它会创建大量的线程来并行执行相同的代码。每个线程都有一个唯一的 ID，可用于区分它与其他线程，并访问全局内存中的不同数据部分。

线程被组织成线程块 (thread blocks)，线程块又被组织成网格 (grids)。这种层级结构允许程序员将复杂问题分解为可以在 GPU 上高效执行的更小、更易于管理的部分。

线程束 (Warp)

线程束是 CUDA 编程模型中 32 个线程的分组。它是 GPU 上的基本执行单位。

一个线程束中的线程在 GPU 核心上以 SIMT（单指令多线程）方式执行。

一个线程束中的所有线程同时执行相同的指令。这种执行模型被称为单指令多线程（SIMT）。如果一个线程束中的线程需要执行不同的指令（例如，由于 if 语句），就会发生线程束分歧 (warp divergence)，这可能会降低性能。

线程束由流式多处理器（SM）上的线程束调度器 (Warp Scheduler)进行调度。

线程束组 (Warpgroup)

线程束组是 CUDA 编程模型中线程束 (warps) 的分组。

在 Hopper 架构中，一个线程束组由四个线程束组成。一个线程束组内的线程束可以在流式多处理器（SM）上协同执行，共享资源并同步它们的执行。

协作线程数组 (Cooperative Thread Array)

协作线程数组（CTA）是 CUDA 编程模型中线程 (threads) 的分组。它与线程块 (thread block) 是同义词。

协作线程数组对应于 CUDA 编程模型中线程块层级的线程块。修改自 NVIDIA 的 CUDA Refresher 和 CUDA C++ 编程指南。

“协作线程数组”这个术语强调了一个线程块内的线程可以相互协作，例如通过共享内存和同步。

核函数 (Kernel)

在 CUDA 中，核函数是一个在 GPU 上并行执行的函数。它由大量线程 (threads) 同时执行，每个线程都在数据集的不同部分上运行相同的代码。

单个核函数启动对应于 CUDA 编程模型中的一个线程块网格。修改自 NVIDIA 的 CUDA Refresher 和 CUDA C++ 编程指南。

核函数是使用 CUDA C++ 编写的，并从主机（CPU）代码中启动。当一个核函数被调用时，它会创建一个线程块网格 (grid of thread blocks)，每个线程块包含一组线程。

线程块 (Thread Block)

线程块是 CUDA 编程模型中线程 (threads) 的分组。它是协作线程数组（Cooperative Thread Array）的同义词。

线程块是 CUDA 编程模型中线程层级结构的一个级别。修改自 NVIDIA 的 CUDA Refresher 和 CUDA C++ 编程指南。

一个线程块内的线程可以相互协作，例如通过共享内存和同步。这使得它们非常适合执行需要数据共享和协调的任务。

线程块网格 (Thread Block Grid)

线程块网格（或简称网格）是 CUDA 编程模型中线程块 (thread blocks) 的分组。

线程块网格是 CUDA 编程模型中线程层级结构的最高级别。修改自 NVIDIA 的 CUDA Refresher 和 CUDA C++ 编程指南。

当一个核函数 (kernel) 启动时，它会创建一个由一个或多个线程块组成的网格。一个网格内的线程块独立执行，不能直接相互通信。

线程层级 (Thread Hierarchy)

CUDA 线程层级是指 CUDA 编程模型中线程的组织结构。

CUDA 编程模型的线程层级结构，从单个线程到线程块，再到线程块网格（左侧），映射到硬件上，从 CUDA 核心到流式多处理器，再到整个 GPU（右侧）。修改自 NVIDIA 的 CUDA Refresher 和 CUDA C++ 编程指南。

该层级结构由三个级别组成：

  1. 线程 (Thread)：执行计算的最小单位。

  2. 线程块 (Thread Block)：一组可以相互协作的线程。

  3. 网格 (Grid)：一组独立执行的线程块。

这种层级结构允许程序员将复杂问题分解为可以在 GPU 上高效执行的更小、更易于管理的部分。

内存层级 (Memory Hierarchy)

CUDA 内存层级是指 CUDA 编程模型中可供线程 (threads) 使用的不同类型的内存。

共享内存和全局内存是 CUDA 编程模型（左）中内存层级的两个级别，分别映射到 L1 数据缓存和 GPU RAM。修改自 NVIDIA 的 CUDA Refresher 和 CUDA C++ 编程指南。

该层级结构由以下几个级别组成：

  • 寄存器 (Registers)：每个线程私有的最快内存。

  • 共享内存 (Shared Memory)：每个线程块内的所有线程共享的内存。

  • 全局内存 (Global Memory)：所有线程均可访问的最慢但最大的内存。

了解内存层级对于编写高性能的 CUDA 核函数 (kernels) 至关重要。通过有效利用不同级别的内存，开发者可以最小化内存延迟 (latency)并最大化带宽 (bandwidth)。

寄存器 (Registers)

寄存器是 CUDA 内存层级中最小、最快的内存空间。

寄存器是与单个线程（左）关联的内存层级的内存。修改自 NVIDIA 的 CUDA Refresher 和 CUDA C++ 编程指南。

每个线程 (thread) 都有一组私有寄存器，用于存储其本地变量和计算的中间结果。寄存器由编译器管理，不能由程序员直接寻址。

对寄存器的访问速度非常快，通常只需要一个时钟周期。然而，寄存器的数量是有限的，过度使用寄存器（称为寄存器压力 (register pressure)）可能会影响性能。

共享内存 (Shared Memory)

共享内存是位于流式多处理器（SM）上的一个高速内存区域，由一个线程块内的所有线程共享。

共享内存是 CUDA 编程模型中内存层级的一个级别。修改自 NVIDIA 的 CUDA Refresher 和 CUDA C++ 编程指南。

共享内存对于需要数据共享和协作的并行算法至关重要。通过在共享内存中暂存数据，一个线程块内的线程可以避免对较慢的全局内存进行冗余访问。

与全局内存相比，共享内存的延迟要低得多，带宽也要高得多。然而，它的容量要小得多，并且其生命周期仅限于其所在线程块的执行期间。

全局内存 (Global Memory)

全局内存是 CUDA 内存层级中最大、但最慢的内存空间。它可以被 GPU 上的所有线程以及主机（CPU）访问。

全局内存是 CUDA 编程模型中内存层级的最高级别。它存储在 GPU RAM 中。修改自 NVIDIA 的 CUDA Refresher 和 CUDA C++ 编程指南。

全局内存对于在主机和设备之间传输数据以及存储无法放入较小、较快内存（如共享内存或寄存器）的大型数据集至关重要。

对全局内存的访问具有高延迟 (latency)和有限的带宽 (bandwidth)。因此，最小化全局内存访问次数对于实现高性能的 CUDA 核函数 (kernels) 至关重要。

合并内存访问 (coalesced memory access)等技术可用于优化全局内存访问并提高性能。

主机软件 (Host Software)

主机软件是指在主机（CPU）上运行并与 GPU 交互的程序和库。

CUDA（软件平台）

CUDA 软件平台是一个全面的工具和库生态系统，使开发者能够利用 NVIDIA GPU 的强大能力进行通用计算。

CUDA 软件平台由多个层次组成，从低级的 GPU 驱动程序和库到高级的编程语言和应用程序。

该平台的核心组件包括：

  • NVIDIA GPU 驱动程序：允许主机与 GPU 通信的低级软件。

  • CUDA C++：一种 C++ 的扩展，允许开发者编写在 GPU 上执行的核函数 (kernels)。

  • CUDA 库：一个针对各种领域（如线性代数 (cuBLAS)、深度学习 (cuDNN) 和信号处理）优化的 GPU 加速库集合。

  • CUDA 开发工具：一系列用于分析、调试和优化 CUDA 应用程序的工具，例如 NVIDIA Nsight Systems 和 nvcc。

CUDA C++（编程语言）

CUDA C++ 是 C++ 编程语言的一个扩展，允许开发者为 NVIDIA GPU 编写程序。它引入了一组新的语言扩展，使开发者能够定义和启动在 GPU 上并行执行的核函数 (kernels)。

CUDA C++ 的一些关键特性包括：

  • __global__ 关键字：用于定义一个可以从主机（CPU）调用并在设备（GPU）上执行的核函数。

  • __device__ 关键字：用于定义一个只能从设备（GPU）调用的函数。

  • __host__ 关键字：用于定义一个只能从主机（CPU）调用的函数。

  • 执行配置语法 (<<<...>>>)：用于指定启动核函数时要使用的线程块 (thread blocks) 和线程 (threads) 的数量。

  • 内置变量：例如 threadIdx、blockIdx、blockDim 和 gridDim，它们提供有关线程层级中每个线程的索引和维度的信息。

NVIDIA GPU 驱动程序 (NVIDIA GPU Drivers)

NVIDIA GPU 驱动程序是一组低级软件组件，允许操作系统和应用程序与 NVIDIA GPU 硬件进行通信。

驱动程序负责管理 GPU 的资源，例如内存、核心和调度器。它们还提供了一组 API，使开发者能够访问 GPU 的功能。

驱动程序由以下几个部分组成：

nvidia.ko

nvidia.ko 是 NVIDIA GPU 驱动程序的内核模式组件。它是一个可加载的内核模块，直接与操作系统内核交互，以管理 GPU 硬件。

nvidia.ko 负责以下任务：

.ko 文件扩展名代表“内核对象（Kernel Object）”。

CUDA 驱动 API (CUDA Driver API)

CUDA 驱动 API 是一个低级编程接口，提供对 NVIDIA GPU 硬件的直接控制。它被 CUDA 运行时 API 和其他高级库用于与 GPU 交互。

驱动 API 允许开发者：

  • 管理 GPU 设备和上下文

  • 分配和释放设备内存

  • 在主机和设备之间传输数据

  • 启动 CUDA 核函数 (kernels)

  • 同步主机和设备的执行

驱动 API 在 libcuda.so 库中实现。

libcuda.so

libcuda.so 是实现 CUDA 驱动 API 的共享库。它负责与 nvidia.ko 内核模式驱动程序通信，以管理 GPU 硬件。

当一个 CUDA 应用程序启动时，它会加载 libcuda.so 库并使用它来与 GPU 交互。.so 文件扩展名代表“共享对象（Shared Object）”。

NVIDIA 管理库 (NVIDIA Management Library, NVML)

NVML 是一个用于监控和管理 NVIDIA GPU 的库。它提供了一个 C 语言 API，允许开发者查询 GPU 的状态，例如：

  • GPU 利用率

  • 内存使用情况

  • 温度

  • 功耗

  • 时钟速度

NVML 还在 nvidia-smi 命令行工具中使用。

libnvml.so

libnvml.so 是实现 NVML API 的共享库。.so 文件扩展名代表“共享对象（Shared Object）”。

nvidia-smi

nvidia-smi（NVIDIA 系统管理接口）是一个命令行工具，用于监控和管理 NVIDIA GPU。它使用 NVML 库来查询 GPU 的状态并显示有关以下内容的信息：

  • GPU 型号和驱动程序版本

  • GPU 利用率和内存使用情况

  • 温度和功耗

  • 正在运行的进程

  • 时钟速度

CUDA 运行时 API (CUDA Runtime API)

CUDA 运行时 API 是一个高级编程接口，简化了 NVIDIA GPU 的编程。它在 CUDA 驱动 API 之上提供了一个抽象层，使开发者更容易编写 CUDA 应用程序。

运行时 API 负责以下任务：

  • GPU 初始化和上下文管理

  • 内存分配和传输

  • 核函数启动和同步

运行时 API 在 libcudart.so 库中实现。

libcudart.so

libcudart.so 是实现 CUDA 运行时 API 的共享库。.so 文件扩展名代表“共享对象（Shared Object）”。

NVIDIA CUDA 编译器驱动 (NVIDIA CUDA Compiler Driver, nvcc)

nvcc 是用于编译 CUDA C++ 代码的编译器驱动。它负责将 CUDA 核函数 (kernels) 编译成可以在 NVIDIA GPU 上执行的代码。

nvcc 编译器执行以下步骤：

  1. 将主机代码（在 CPU 上运行）与设备代码（在 GPU 上运行）分离。

  2. 将设备代码编译成 PTX，这是一种中间表示。

  3. 将主机代码编译成目标 CPU 架构的目标文件。

  4. 将 PTX 代码嵌入到主机目标文件中。

  5. 将主机目标文件链接到可执行文件中。

NVIDIA 运行时编译器 (NVIDIA Runtime Compiler, NVRTC)

NVRTC 是一个在运行时将 CUDA C++ 代码编译成 PTX 的库。它允许应用程序在运行时动态生成和编译 CUDA 核函数 (kernels)，而无需使用 nvcc 编译器。

NVRTC 对于需要根据用户输入或运行时条件生成专用核函数的应用程序非常有用。

NVIDIA CUDA 分析工具接口 (NVIDIA CUDA Profiling Tools Interface, CUPTI)

CUPTI 是一个用于分析和追踪 CUDA 应用程序的 API。它允许开发者收集有关 CUDA 核函数 (kernels) 执行、内存传输和 API 调用的详细信息。

CUPTI 被 NVIDIA Nsight Systems 等分析工具用于收集数据和可视化应用程序的性能。

NVIDIA Nsight Systems

NVIDIA Nsight Systems 是一个性能分析工具，可帮助开发者理解和优化其应用程序的行为。它收集并显示有关 CPU 和 GPU 活动的详细信息，包括：

Nsight Systems 使用 CUPTI 库来收集 GPU 活动数据。

CUDA 二进制实用程序 (CUDA Binary Utilities)

CUDA 二进制实用程序是一组用于处理 CUDA 二进制文件的工具，例如：

  • cuobjdump：显示有关 CUDA 二进制文件的信息，包括核函数 (kernels)、PTX 和 SASS 代码。

  • nvdisasm：反汇编 CUDA 二进制文件并显示 SASS 代码。

cuBLAS

cuBLAS 是一个由 NVIDIA 提供的 GPU 加速的基本线性代数子程序（BLAS）库。它为常见的矩阵和向量运算提供了高度优化的实现，例如：

  • 矩阵-向量乘法 (GEMV)

  • 矩阵-矩阵乘法 (GEMM)

  • 求解三角系统 (TRSM)

cuBLAS 对于加速在科学计算、工程和机器学习等领域严重依赖线性代数运算的应用程序至关重要。

cuDNN

cuDNN（CUDA 深度神经网络库）是一个由 NVIDIA 提供的 GPU 加速的深度神经网络基元库。它为深度学习中常见的例程提供了高度优化的实现，例如：

  • 卷积

  • 池化

  • 归一化

  • 激活函数

cuDNN 被 TensorFlow、PyTorch 和 MXNet 等深度学习框架广泛用于加速神经网络的训练和推理。

性能 (Performance)

GPU 性能是指 GPU 执行计算任务的速度和效率。

性能瓶颈 (Performance Bottleneck)

性能瓶颈是系统中限制整体性能的一个组件。在 GPU 计算中，瓶颈可能是：

像这样的 Roofline 图用于快速识别吞吐量导向系统中的性能瓶颈。改编自 Williams, Waterman, and Patterson (2008)。

屋顶线模型 (Roofline Model)

屋顶线模型是一种直观的性能模型，用于确定计算核函数 (kernel) 是受计算限制还是受内存限制。

在屋顶线模型中，运算/算术强度绘制在水平轴上。图表改编自 Williams, Waterman, and Patterson (2008)。

该模型根据 GPU 的峰值算术带宽（以 FLOPS/秒为单位）和峰值内存带宽（以字节/秒为单位），绘制了可实现的性能与算术强度（以 FLOPS/字节为单位）的关系图。

“屋顶线”代表了 GPU 的理论峰值性能。位于屋顶线平坦部分的核函数被认为是计算限制的，而位于倾斜部分的核函数则被认为是内存限制的。

计算限制 (Compute-bound)

当 GPU 应用程序的性能受其执行计算的速度限制，而不是受其访问内存的速度限制时，该程序就是计算限制的。

在屋顶线图中，计算限制的程序位于图表的平坦部分。图表改编自 Williams, Waterman, and Patterson (2008)。

换句话说，GPU 花在执行算术运算上的时间比花在等待数据从内存中获取的时间要多。

具有高算术强度（即每次内存访问执行大量计算）的程序更有可能是计算限制的。

内存限制 (Memory-bound)

当 GPU 应用程序的性能受其访问内存的速度限制，而不是受其执行计算的速度限制时，该程序就是内存限制的。

像上面这样的屋顶线图有助于确定程序的性能瓶颈是计算能力、内存带宽还是其他因素。图表改编自 Williams, Waterman, and Patterson (2008)。

换句话说，GPU 花在等待数据从内存中获取的时间比花在执行算术运算上的时间要多。

具有低算术强度（即每次内存访问执行少量计算）的程序更有可能是内存限制的。

算术强度 (Arithmetic Intensity)

算术强度（有时称为运算强度）是衡量一个计算核函数 (kernel) 在每个数据字节上执行的算术运算量的指标。它通常以 FLOPS/字节（每字节浮点运算次数）为单位进行测量。

在屋顶线模型中，运算/算术强度绘制在水平轴上。图表改编自 Williams, Waterman, and Patterson (2008)。

高算术强度的核函数对每个加载到 GPU 的数据元素执行大量计算。这些核函数通常是计算限制的。

低算术强度的核函数对每个加载到 GPU 的数据元素执行少量计算。这些核函数通常是内存限制的。

开销 (Overhead)

在并行计算中，开销是指不直接有助于最终结果的任何计算或资源。

在 GPU 编程中，开销可能来自多种来源，包括：

最小化开销对于实现高性能的 CUDA 应用程序至关重要。

利特尔定律 (Little's Law)

利特尔定律是排队论中的一个基本定理，它描述了系统中项目的平均数量（L）、项目的平均到达率（λ）和项目在系统中花费的平均时间（W）之间的关系。

该定律通常表示为：L = λ * W

在 GPU 计算中，利特尔定律可用于推断系统的吞吐量 (throughput)、延迟 (latency) 和并发性。例如，如果一个 GPU 可以同时处理 16,000 个线程 (threads)，每个线程需要 100 微秒完成，那么 GPU 的吞吐量为：

16,000 / 100 µs = 160,000,000 线程/秒

内存带宽 (Memory Bandwidth)

内存带宽是 GPU 内存与流式多处理器（SM）之间数据传输的速率。它通常以千兆字节/秒（GB/s）为单位进行测量。

内存带宽是 GPU 性能的关键因素，特别是对于内存限制的应用程序。高内存带宽允许 GPU 快速地将数据输入和输出其核心，从而减少等待数据的时间。

算术带宽 (Arithmetic Bandwidth)

算术带宽（也称为计算吞吐量）是 GPU 执行算术运算的速率。它通常以每秒浮点运算次数（FLOPS）或每秒整数运算次数（IOPS）来衡量。

算术带宽是 GPU 性能的关键因素，特别是对于计算限制的应用程序。

延迟隐藏 (Latency Hiding)

延迟隐藏是一种通过在等待长延迟操作（如内存访问）完成时执行其他工作来重叠计算和内存传输的技术。

在 GPU 计算中，延迟隐藏是通过利用大量的并发 (concurrent) 线程 (threads) 来实现的。当一个线程束 (warp) 因内存访问而停滞时，线程束调度器 (Warp Scheduler) 可以切换到另一个准备好执行的线程束，从而保持 GPU 的核心繁忙。

线程束执行状态 (Warp Execution State)

线程束执行状态描述了线程束 (warp) 在流式多处理器（SM）上的当前状态。

线程束可以处于以下几种状态之一：

  • active：线程束正在 SM 上积极执行。

  • eligible：线程束已准备好执行，但正在等待 SM 上的可用资源。

  • stalled：线程束正在等待长延迟操作（如内存访问或同步）完成。

了解线程束执行状态对于识别性能瓶颈和优化 CUDA 核函数 (kernels) 非常重要。

活动周期 (Active Cycle)

活动周期是指流式多处理器（SM）中至少有一个线程束 (warp) 正在执行指令的时钟周期。

此图中描绘的所有周期都是活动周期。图表灵感来自 GTC 2025 的 CUDA Techniques to Maximize Compute and Instruction Throughput 演讲。

活动周期的数量可以用来衡量 SM 的利用率。高比例的活动周期表示 SM 得到了很好的利用，而低比例的活动周期则可能表示性能瓶颈，例如内存延迟 (latency)或线程束分歧 (warp divergence)。

占用率 (Occupancy)

占用率是衡量流式多处理器（SM）上活动线程束 (warps) 数量与 SM 支持的最大活动线程束数量之比的指标。

高占用率对于隐藏延迟 (latency hiding) 和最大化 GPU 利用率非常重要。但是，高占用率并不总能保证高性能。其他因素，如内存合并 (memory coalescing)和线程束分歧 (warp divergence)，也会影响性能。

流水线利用率 (Pipe Utilization)

流水线利用率是衡量 GPU 流式多处理器（SM）内部特定计算流水线繁忙程度的指标。

SM 包含多个专门的流水线，用于不同类型的操作，例如：

  • 整数和浮点运算

  • 内存加载和存储

  • 特殊函数（例如 sin、cos）

高流水线利用率表示特定的计算资源得到了很好的利用，而低利用率则可能表明瓶颈在于其他地方，或者工作负载没有充分利用可用的硬件。

峰值速率 (Peak Rate)

峰值速率是指 GPU 核心可以执行操作的最大速率。它通常以每周期操作数来衡量。

例如，一个 CUDA 核心可能能够以每周期一个 FMA（融合乘加）操作的峰值速率执行 FP32 运算。

将 GPU 的实际性能与其峰值速率进行比较，可以帮助识别性能瓶颈和优化机会。

发出效率 (Issue Efficiency)

发出效率是衡量在活动周期内实际发出了指令的线程束 (warps) 的比例。

在此图的四个时钟周期中，有三个周期发出了指令，因此发出效率为 75%。图表灵感来自 GTC 2025 的 CUDA Techniques to Maximize Compute and Instruction Throughput 演讲。

例如，如果一个流式多处理器（SM）在 100 个活动周期内为 80 个线程束发出了指令，那么发出效率为 80%。

低发出效率可能表明 SM 没有得到充分利用，这可能是由于线程束分歧 (warp divergence)、内存延迟 (latency)或占用率 (occupancy) 不足等因素造成的。

流式多处理器利用率 (Streaming Multiprocessor Utilization)

SM 利用率是衡量流式多处理器（SM）在 CUDA 核函数 (kernel) 执行期间繁忙程度的指标。它通常表示为 SM 处于活动状态的时钟周期百分比。

高 SM 利用率表示 GPU 的计算资源得到了很好的利用。低 SM 利用率则可能表明瓶颈在于其他地方，例如内存带宽或 PCI-e 传输。

线程束分歧 (Warp Divergence)

当一个线程束 (warp) 内的线程 (threads) 采取不同的执行路径时（例如，由于 if-else 语句），就会发生线程束分歧。

由于一个线程束内的所有线程必须执行相同的指令，因此分歧的线程会被序列化，这意味着它们会依次执行不同的代码路径。这可能会导致显著的性能下降，因为一些线程在等待其他线程完成其路径时处于空闲状态。

记分板停顿 (Scoreboard Stall)

记分板停顿是一种由于数据依赖性而发生的停顿。当一条指令需要等待前一条指令完成才能执行时，就会发生这种情况。

在 GPU 中，记分板用于跟踪寄存器的可用性。当一条指令需要一个正在被前一条指令使用的寄存器时，记分板会使该指令停顿，直到该寄存器可用为止。

分支效率 (Branch Efficiency)

分支效率是衡量 GPU 处理分支指令（例如 if-else 语句）效率的指标。

由于线程束分歧 (warp divergence)，分支指令可能会在 GPU 上产生高昂的代价。当一个线程束 (warp) 内的线程 (threads) 采取不同的分支路径时，这些路径会被序列化，从而导致性能下降。

高分支效率表示大多数线程都采取相同的分支路径，从而最小化了线程束分歧。

合并内存访问 (Coalesced Memory Access)

合并内存访问是一种内存访问模式，其中一个线程束 (warp) 内的所有线程 (threads) 同时访问全局内存中的连续位置。

当内存访问被合并时，GPU 可以在一次内存事务中满足所有线程的请求。这大大减少了内存事务的总数，并提高了内存带宽。

非合并访问会导致多次内存事务，从而导致性能下降。

存储体冲突 (Bank Conflict)

当多个线程 (threads) 同时尝试访问共享内存的同一个存储体时，就会发生存储体冲突。

在此图中，线程 0 和线程 2 都访问了存储体 0，导致了存储体冲突。

共享内存被划分为多个大小相等的存储体。当多个线程访问不同的存储体时，可以并行访问。但是，如果多个线程访问同一个存储体，访问将被序列化，从而导致性能下降。

寄存器压力 (Register Pressure)

当一个 CUDA 核函数 (kernel) 需要的寄存器数量超过了流式多处理器（SM）上可用的寄存器数量时，就会发生寄存器压力。

当发生这种情况时，编译器必须将一些变量“溢出”到较慢的本地内存中。这可能会导致显著的性能下降，因为对本地内存的访问比对寄存器的访问要慢得多。

寄存器压力可能会限制 SM 上的占用率 (occupancy)，因为 SM 需要在活动线程束 (warps) 之间划分可用的寄存器。

贡献者 (Contributors)

我们衷心感谢以下为本术语表做出贡献的个人：

技术文章
GPU
License:  CC BY 4.0