在上一篇文章《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 在科學計算應用中的實用性。

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 GFLOPS

2.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