Fermi (微架構)

費米（）是英伟达在2010年3月发布的圖形處理器微架构的代号，是架构的继承者，于2010年4月正式推向零售市场。

費米是英伟达最後一款在遊戲顯卡上保留完整雙精度的微架构。大多数GeForce 400 系列、GeForce 500 系列和部分GeForce 800 系列显卡基于費米架构，桌面版Fermi GPU均采用40nm工艺制造，而移动版则采用了40nm和28nm两种工艺制程。在工作站市场方面，費米架构也应用于NVIDIA Quadro x000系列、Quadro NVS型号以及NVIDIA Tesla计算模块系列。Fermi是英伟达截止到目前最古老的、仍支持Microsoft Direct3D 12（feature_level_11）渲染API的微架构。

費米的后继者为开普勒，但其仍在GeForce 600系列、700系列和800系列显卡中曾与Kepler一同使用（后两者仅限移动版 GPU）。

費米架构的命名来自20世纪的美籍義大利裔物理學家恩里科·費米（）。

概述
NVIDIA GeForce GTX 480是GeForce 400系列显卡中的首款产品，采用了Fermi微架构（核心代号 GF100-375-A3）。其主要规格如下：

• 流式多处理器 (SM)：每个 SM 包含 32 个CUDA核心（详见“流式多处理器”和"CUDA 核心”部分）。
• GigaThread 全局调度器：负责将线程块分发给SM的线程调度器，并管理执行过程中的线程上下文切换（详见"Warp 调度”部分）。
• 主机接口：通过PCI-Express 2.0总线连接GPU与CPU，峰值传输速率为8 GB/s。
• 显存 (DRAM)：得益于 64 位寻址能力，支持高达6 GB的GDDR5显存（详见“内存架构”部分）。
• 时钟频率：1.5 GHz（NVIDIA 未正式公布此数据，由 Insight 64 估算得出）。
• 峰值性能：1.5 TFlops。
• 全局内存时钟：2 GHz。
• 显存带宽：192 GB/s。
• 视频解码支持：支持 H.264 全高清 (FHD) 解码。

流式多处理器 (Streaming Multiprocessor)
每个 SM 单元包含以下组件：

• 32 个单精度 CUDA 核心；
• 16 个加载/存储单元；
• 4 个特殊功能单元 (SFUs)；
• 64 KB 片上高速内存块；
• L2 缓存接口。

加载/存储单元 (Load/Store Units)

• 每个时钟周期可为16个线程计算源地址和目的地址。
• 负责从缓存或DRAM中加载数据或将数据存储到这些区域。

特殊功能单元 (SFUs)

• 执行超越函数指令，例如正弦 (sin)、余弦 (cosine)、倒数 (reciprocal) 和平方根 (square root)。
• 每个 SFU 在每个时钟周期为每个线程执行一条指令；一个 Warp（线程束）完成这些操作大约需要八个时钟周期。
• SFU 流水线与分发单元解耦，这意味着即使 SFU 正在被占用，分发单元也可以向其他执行单元发出指令，从而提高整体效率。

CUDA 核心 (CUDA Core)

• 整数算术逻辑单元 (ALU)

** 支持所有指令的完整32位精度，符合标准编程语言的要求。
** 针对工作站和服务器型号进行了优化，以高效支持 64 位运算；但在消费级版本中，该功能被人为限制（ crippled）。

• 浮点运算单元 (FPU)

** 实现了新的 IEEE 754-2008 浮点数标准。
** 为单精度和双精度算术均提供了融合乘加 (FMA) 指令。
** 每个 SM 在每个时钟周期最多可执行 16 次双精度融合乘加运算。

• 融合乘加 (Fused Multiply-Add, FMA)

** 将乘法与加法（即 A×B+C）合并为一次操作，仅进行一次最终的舍入步骤。
** 在加法过程中不会丢失精度。
** 相比于分开执行乘法和加法运算，FMA 具有更高的计算准确性。

Warp 调度 (Warp Scheduling)
Fermi 架构采用两级分布式线程调度器。

• 指令发射机制：每个 SM（流式多处理器）可以并发发射指令，利用四个执行列中的任意两列。例如，SM 可以将来自第一列核心的 16 个操作与第二列核心的 16 个操作混合执行；或者将来自加载/存储单元的 16 个操作与来自 SFU（特殊功能单元）的 4 个操作组合；也可以根据程序需求进行任何其他组合。

注意**：64 位浮点运算需要同时使用前两列执行单元，因此其运行速度仅为 32 位运算的一半。

• 双 Warp 调度器 (Dual Warp Scheduler)：

** 在 SM 层面，每个 Warp 调度器负责将包含 32 个线程的 Warp 分发给其执行单元。
** 每个 SM 配备两个 Warp 调度器和两个指令分发单元，这使得 SM 能够同时发射并执行两个 Warp。
双发射能力**：双 Warp 调度器会选择两个 Warp，并从每个 Warp 中各取一条指令，分发给一组 16 个核心、16 个加载/存储单元或 4 个 SFU。大多数指令支持“双发射”（Dual Issued），即可以并发执行两条整数指令、两条浮点指令，或者混合执行整数、浮点、加载、存储及 SFU 指令。
限制**：双精度（64 位）指令不支持与其他任何操作进行双分发。

性能 (Performance)

• 单精度理论算力 (Single-Precision GFLOPS)：

** Fermi GPU 的单精度理论峰值性能计算公式为：2 × CUDA 核心数量 × Shader 时钟频率 (GHz)。
*** 其中系数 2 代表每个 FMA（融合乘加）指令在每个时钟周期内完成的操作次数。
架构限制说明**：上一代 Tesla 架构允许在 CUDA 核心和 SFU（特殊功能单元）之间并行双发射 MAD+MUL 指令，但 Fermi 架构失去了这一能力。Fermi 每个 SM 每周期仅能发射 32 条指令，这刚好使其 32 个 CUDA 核心保持满负荷运行，而无法利用额外的 SFU 资源来进一步提升单核周期的操作数（即无法超过每个 CUDA 核心每周期 2 次操作）。

• 双精度理论算力 (Double-Precision GFLOPS)：

Fermi GPU 的双精度理论性能约为其单精度性能的 1/2**（针对 GF100/GF110 架构）。
实际可用性限制**：
专业卡*：双精度算力仅在专业的 Quadro 和 Tesla 系列显卡上完全可用。
消费级卡：GeForce 等消费级显卡的双精度性能被人为限制（artificially capped）为单精度性能的 1/8*。

内存系统 (Memory)

• L1 缓存与 L2 缓存：

L1 Cache**：每个 SM（流式多处理器）拥有独立的 L1 缓存。
Unified L2 Cache**：统一的 L2 缓存服务于所有操作，包括加载、存储和纹理请求。

• 寄存器 (Registers)：

容量与隔离性**：每个 SM 拥有 32K（即 32,768）个 32 位寄存器。线程仅能访问自己的寄存器，无法访问其他线程的寄存器。
使用限制**：CUDA 内核最多可使用 63 个寄存器。随着线程数增加（工作负载加重），可用寄存器数量会从 63 平滑下降至 21。
带宽**：寄存器具有极高的内部带宽，约为 8,000GB/s。

• L1 + 共享内存 (L1+Shared Memory)：

功能**：片上内存，可配置为两种模式：
**# 48 KB 共享内存 + 16 KB L1 缓存；
**# 16 KB 共享内存 + 48 KB L1 缓存。
用途**：
L1 Cache*：用于为单个线程缓存数据（寄存器溢出/Spilling）。
Shared Memory*：用于同一线程块内的多个线程间共享数据。它支持线程协作、促进片上数据的重用，并大幅减少片外内存流量。
特性**：由同一线程块内的线程访问，提供低延迟（10-20 个时钟周期）和高带宽（1,600 GB/s），适用于中等规模数据（如计算中间结果、矩阵运算的行/列、视频行等）。David Patterson 指出这种共享内存采用了“局部 Scratchpad"的设计理念。

• 本地内存 (Local Memory)：

定义**：用于存放“溢出”寄存器的内存位置。当线程块所需的寄存器存储空间超过 SM 可用容量时，会发生寄存器溢出（Register Spilling）。
使用场景**：仅用于某些自动变量（在设备代码中声明，且未加 __device__、__shared__ 或 __constant__ 限定符的变量）。
分配规则**：通常变量驻留在寄存器中，但在以下情况会被分配到本地内存：
**# 编译器无法确定数组索引为常量的数组；
**# 占用寄存器空间过大的大型结构体或数组；
**# 内核使用的寄存器总数超过 SM 可用数量时，编译器决定溢出的变量。

• L2 Cache：

规格**：768 KB 统一 L2 缓存，由 16 个 SM 共享。
功能**：服务于所有来自/去往全局内存的加载和存储操作（包括与 CPU 主机的数据拷贝），以及纹理请求。
原子操作**：L2 缓存子系统还实现了原子操作，用于管理需要在线程块甚至内核间共享的数据访问。

• 全局内存 (Global Memory / VRAM)：

访问方式**：所有线程和主机系统（通过 PCIe 总线）均可直接访问。
延迟**：具有较高延迟，约为 400-800 个时钟周期。

视频编解码 (Video Decompression/Compression)

• 硬件支持现状：

该架构（Fermi）集成了 Nvidia NVDEC（前身为 NVCUVID）硬件单元，用于加速视频的解压缩**。
它也继承了上一代技术 PureVideo** 的部分功能。

• 编码能力缺失：

Fermi 架构不支持视频编码**（Compression）。
专用的视频编码技术 Nvidia NVENC 尚未在该架构中引入，而是直到其继任者 Kepler** 架构才首次出现。