聊聊英伟达于2010年发布的第一个完整 GPU 计算架构！
张彦飞allen约 8 分钟大家好，我是飞哥！
在上一篇文章《NVIDIA通用计算首代架构 Tesla 与 CUDA 1.0 剖析》中我们提到了，首代 Tesla G80 通过统一着色架构（Unified Shader Architecture）将传统分离的顶点着色器、像素着色器、几何着色器统一为可编程的流处理器 SP（Streaming Processor）。SP 虽然设计上主要是处理可执行顶点、像素等着色图形任务，但也可执行与图形无关的计算指令，为通用计算提供了非常好的基础。
虽然 Tesla G80 后面经过了一系列的升级，包括 G92、GT200 等核心。但是整体上来，Tesla 架构还是不能很好地满足科学计算的需求。最大的制约是对于双精度浮点数 FP64 的运算性能太差了。即便是在 GT200 核心中，一个时钟周期也只能进行 30 次 的双精度乘加 FMA（Fused Multiply-Add operation）。
所以，英伟达在 2010 年又推出了 Fermi 架构。在推出后，英伟达把它称为第一个完整的 GPU 计算架构。我们今天进来看看 Fermi 都带来了哪些新的改进。为什么英伟达把它称为第一个完整的 GPU架构。
一、Fermi 架构提升
1.1 制程工艺
Fermi 架构中的核心参数首先值得提的是它的制程工艺。
在 Tesla 架构的 G80 核心中采用的是台积电（TSMC） 的 90 nm 工艺。由于工艺的限制，散热就是一个大问题。所以 Tesla GPU 中的晶体管数量就会受到限制。G80 芯片中总共包含了大约 7 亿个晶体管。
到了 2008 年 Tesla 架构下的 GT200 核心，制程工艺有了较大进步，使用了 55 nm 工艺。晶体管数量提升到了 14 亿。从 SP 流处理器（2010 年后开始称CUDA 核心）数量上来看上，从 128 -> 240 个。
到了 2010 年 Fermi 架构下的 GF100 开始采用 40 nm 的工艺。因为制程工艺的提升，芯片晶体管数量和SM数量、CUDA核心数都有大幅度提升。单芯片晶体管数量提升到了 30 亿。CUDA 核心数更是提升到了 512 个。
可见，制程工艺的提升对于芯片的提升帮助是非常大的。
1.2 显存系统
另外就是 Fermi 在显存上也有较大的提升。
从显存代际上来看，Fermi 升级到了 GDDR5 代际（G80使用的是GDDR3）。更高的显存代际意味着显存的数据频率就会越高。G80 GeForce 8800 Ultra 中采用的 GDDR3 的有效数据频率为2.2 Gbps。而 Tesla M2070 的 有效数据频率为可达 3.1 Gbps effective。
另外就是对 ECC 的支持。在科学计算应用场景中，服务器经常是部署在数据中心中的。在数据中心中对数据错误的容忍度很低。所以对 ECC 内存的支持也大幅提升了 Fermi 在科学计算应用中的实用性。关于ECC内存原理参考飞哥之前的文章服务器之 ECC 内存的工作原理
1.3 缓存系统
在 CPU 中，我们早已经习惯了在 CPU 硬件内部包含 L1、L2、L3 等缓存，用来加速对内存中数据的访问。同样地，这个性能优化思路也可以用在 GPU 中。
在第一代 Tesla G80 核心中是没有设计 L2 缓存系统的。SM 需要访问的数据直接通过 DRAM 来读取。
到了 Fermi 架构下的 GF100 核心，在 Die 的中央位置设计了一块 768 KB 的 L2 缓存。有了这块 L2 缓存，GPU 可将数据一次性从显存加载后供所有 SM 复用，减少对显存的重复访问，降低显存带宽压力。
另外就是 L1 缓存。在 Tesla G80 的每个 SM 内有一个 “共享内存” 组件。程序员需手动通过 __shared__关键字定义变量，并编写数据加载 / 存储代码来管理其数据。如果程序员没有将数据预加载到共享内存，需直接访问全局内存。
而在 Fermi GF100 中这个模块升级到了 64 KB。而且还支持通过配置，分一部分作为 L1 缓存来使用。这样对于 L1 缓存，就可以由硬件自动缓存频繁访问的数据。减少对 DRAM 的访问，提升芯片整体性能。
配置方式包括下面两种方案。
48KB 共享内存 + 16KB L1 缓存 16KB 共享内存 + 48KB L1 缓存
这样，Germi GPU 的存储体系就完善了。
1.4 运算单元
在 Tesla G80 的 SP 中包含了一个 FP Unit。该单元实现了对MAD（Multiply-Add）指令的支持。该指令在完成 乘加运算 a×b+c 时，需两个独立的功能单元，计算时会有精度的损失。
而 Fermi GF100 中的 CUDA Core 中的 FP Unit 实现了对更高效的 FMA（Fused Multiply-Add）指令的支持。该指令在计算乘加运算时，由单个硬件单元完成，计算的效率更高，且没有精度损失。
另外 FMA 还支持双精度，一个 CUDA Core 可以在 1 个时钟周期内完成一次 FP64 运算。Fermi GF100 基于 CUDA Core 中对 FMA 的支持，大幅度提升了科学计算中常用的 FP64 计算能力。这也是 GF100 中很有价值的提升点。双精度计算能力提高后，使得 GPU 在科学计算领域实用性大大提升。
1.5 调度单元
Warp Scheduler 主要负责维护多个 Warp 的状态，监控这些 Warp 中线程的执行情况，根据一定的调度算法，从众多可执行的 Warp 中选择出合适的 Warp，确定哪些 Warp 可以在当前时钟周期内执行。
Dispatch Unit 则负责将 Warp Scheduler 选择出来的 Warp 中的指令，分派给具体的执行单元，如 CUDA 核心、加载 / 存储单元（LD/ST Units）或特殊功能单元（SFU）等，让指令能够在这些执行单元上实际运行。
在每个 SM 中有两个 Warp Scheduler。每个 Warp Scheduler 对应一个 Dispatch Unit（后续的 Kepler 架构对应 2 个 Dispatch Unit）。这样每个 SM 能够并发执行两个 Warp，大大提升了并行执行效率。
例如，一个 Warp Scheduler 选择了一个 Warp，并将其指令信息传递给对应的 Dispatch Unit，该 Dispatch Unit 就会将指令分派到相应的执行单元去执行。与此同时，另一个 Warp Scheduler 和 Dispatch Unit 也在执行类似的操作，从而实现并行处理。
二、Tesla M2070 算力
我们本节以基于 Fermi 架构 GF100 Tesla M2070 来看下它的算力情况
2.1 FP32 算力
在 Fermi 架构下将前一代的 SP 修改为 CUDA Core 后，每个 CUDA Core 每个周期执行一次 FMA（Fused Multiply-Add 融合乘加）运算，包含两次 FP32 操作。
对于 Tesla M2070 来说，其 FP32 算力计算公式如下：
FP32 算力 = 核工作频率 × CUDA核数量 × 每个CUDA核每周期 FP32 操作根据 techpowerup 的数据显示（参考：https://www.techpowerup.com/gpu-specs/tesla-m2070.c1535）其 Shader Clock 频率是 1150 MHz。
在 SM 数量上，虽然 GF100 架构最高可有 16 个 SM，但 M2070 上只启用了 14 个。所以 M2070 实际可用的 CUDA 核数 = 14 个 SM * 每个 SM 32 核 = 448 个。
FP32 算力 = 核工作频率 × SM 数量 × 每个CUDA核每周期 FP32 操作
= 1.15 GHz * 448 个 * 2
= 1,030.4 GFLOPS可以看到，相对上一代基于 Tesla G80 的 GeForce 8800 Ultra 这款 GPU 的 387.1 GFLOPS 来说，M2070 的 FP32 大概是它的 3 倍。
2.2 FP64 算力
对于 Fermi 架构来说，相比较前一代的 GPU 很有优势的地方就是在于新版的 CUDA Core 支持 FP64 运算。每个 CUDA Core 每个时钟周期可以完成一次 FP64 运算。
FP64 算力 = 核工作频率 × SM 数量 × 每个CUDA核每周期 FP64 操作
= 1.15 GHz * 448 个 * 1
= 515.2 GFLOPS2.3 内存带宽
为了方便对比，我们先计算一下 G80 的内存带宽。还是以 GeForce 8800 Ultra 为例，该 GPU 采用的 GDDR3 的有效数据频率为2.2 Gbps。算得内存带宽为 105.6 GB/s。
内存带宽 = 内存位宽 * 数据频率 / 8(换算成字节数)
= 384 bit * 2.2 Gbps / 8
= 105.6 GB/s我们再来看 Fermi 架构下的 Tesla M2070。由于此代际升级 GDDR3 -> GDDR5，所以有效数据频率得到了提升，从2.2 Gbps 提升到了 3.1 Gbps。那么可算的 Tesla M2070 的内存带宽为 148.8 GB/s。
内存带宽 = 内存位宽 * 数据频率 / 8(换算成字节数)
= 384 bit * 3.1 Gbps / 8
= 148.8 GB/s三、CUDA优化
前面我们提到过在 CUDA 1.0 中，程序员需要通过 cudaMalloc 和 cudaFree 申请和释放 GPU 全局内存，还需要借助 cudaMemcpy 在 CPU 与 GPU 之间互相拷贝的方式传输数据。
而到了 Fermi 架构时代，CUDA 也进行了升级。其中最重要的优化是统一了CPU 和 GPU 的虚拟地址空间。开发者通过cudaMallocManaged()分配的内存可被 CPU 和 GPU 任一设备访问，再也不需要手动拷贝数据了。
float *data;
cudaMallocManaged(&data, N * sizeof(float)); // 统一内存分配
// CPU初始化数据
for (int i = 0; i < N; i++) data[i] = i;
// GPU直接访问数据
kernel<<<blocks, threads>>>(data);
cudaDeviceSynchronize(); // 隐式数据传输总结
在科学计算领域，是否支持双精度、是否支持 ECC、是否支持CPU、GPU内存统一寻址是非常重要的需求。这些在 Tesla 时代的 GPU 中都没有解决。
到了 Fermi 架构下，英伟达通过硬件架构 + CUDA 优化解决了以上痛点问题。把这些特性首次整合至单一架构中。
双精度计算提速：大幅度提升 FP64 计算性能，首次使GPU在双精度计算领域具备与传统CPU竞争的能力。 ECC内存校验：首次在消费级GPU中引入硬件级ECC（错误检查与纠正），支持数据中心级的可靠性要求（如Tesla M2050系列），避免因内存错误导致的计算误差，这是传统GPU（如Tesla架构）完全缺失的特性。 统一内存寻址：通过CUDA 3.0支持CPU与GPU的统一地址空间（Unified Addressing），消除了数据在主机与设备间的显式拷贝需求，大幅简化编程模型，而前代架构需手动管理内存传输。 L2 缓存组件：通过引入 L2 缓存组件，大大加速了对 GPU 显存的数据访问。
所以，英伟达认为 Fermi 算是自己第一个完整的 GPU 计算架构，参见 NVIDIA’s Fermi: The First Complete GPU Computing Architecture
最后，求赞、求再看、求转发！