NVIDIA通用计算首代架构 Tesla 与 CUDA 1.0 剖析-51CTO.COM

我们大家都知道英伟达早期是做显卡的，在我上大学的时候，同学们给自己的电脑上配一块英伟达的显卡一般都是为了玩游戏更爽。

但是近些年来随着比特币、大模型等应用场景的火热，英伟达也早就完成了从显卡到通用计算 GPGPU 的华丽转身。造就了今天将近4.3万亿美金市值的巨无霸。

我们今天就从英伟达的第一代数据中心 GPU - Tesla 说起。通过深入地了解第一代 Tesla 架构以及和它配套的 CUDA 1.0。学习完本文后，我们将对以下问题有比较深入的理解。

英伟达为啥要从游戏显卡转型科学计算？
Tesla 中的统一着色架构到底是什么，为什么重要？
通用GPU中 SM 流多处理器都包含了哪些硬件模块？
为什么合理使用常量能大幅度提升 GPU 的计算能力？
CUDA如何实现使用C/C++就能进行GPU编程的？

好了，让我们开始今天的 GPU 学习之旅！

一、英伟达游戏卡的困局

先让飞哥带大家把地球时间进度条倒回到 2004 年。

当时英伟达与 ATI（后被 AMD 收购）的显卡大战进入白热化阶段。尽管英伟达通过 GeForce 6/7 系列占据市场优势，但图形市场增长逐渐放缓，且 AMD 凭借 Radeon 9700 Pro 等产品在部分领域实现反超。此时，英伟达继续找到第二增长曲线。

当时就有这么一个领域，科学领域面临计算瓶颈。传统 CPU 集群成本高且不说，而且且性能低下。部分科学研究者开始利用 GPU 的图形 API 和各种 Shader语言（HLSL、GLSL、Cg等）实现了一些图像分割、CT图像重建、快速傅立叶变换、图像以及音视频编解码等算法。

这就让英伟达眼前一亮，原来 GPU 还可以这么玩。于是灵感就来了，如果能在计算场景提供专用硬件，和更方便使用的编程工具，那是不是就能占据通用计算这块市场。这第二曲线不就打开了么！

说干就干。英伟达在 2006 年推出了 Tesla 系列的 GPU，以及 CUDA 编程套件。开始了从显卡到通用计算的转型。

二、Tesla 架构介绍

英伟达在 2006 年开辟了新的通用计算系列产品名字叫 Tesla。

首代 Tesla 通过统一着色架构（Unified Shader Architecture）将传统分离的顶点着色器、像素着色器、几何着色器统一为可编程的流处理器（Streaming Processor）。该流处理器虽然设计上主要是处理可执行顶点、像素等着色图形任务，但也可执行与图形无关的计算指令，为通用计算提供了基础。

自此后面的英伟达的通用计算 GPUs 就都开始以 Tesla 来命名。例如后面 2016 年的 Tesla P100、2017 年的 Tesla V100、2018 年的 Tesla V4 。

但如果你最近有听过老黄的发布会的话，他嘴里一直在说 Data Center GPUs，而不说 Tesla。原因是在 2018 年发布 Tesla V4 之后，因为特斯拉汽车名声太大了，老黄给就他的通用计算 GPU 改了个名字，叫数据中心 Data Center GPUs 。

另外和 CPU 一样，GPU 中也有架构的概念。第一代的数据中心 GPU 的系列产品名字为 Tesla，架构名字也被命名为了 Tesla。

后面大概每隔两年会有一代新的架构名字，例如 2016 年的 Tesla P100 属于 Pascal 架构。我们熟悉的 2020 年的 A100 ，其架构名是 Amphere。2022年的 H100、H200 以及阉割版的 H20，其架构名是 Hopper。

图片

在这里我们要注意区分 GPU 中的几个重要概念。

第一个是架构名。

架构名代表是一代技术平台，如Tesla（2006、）Maxwell（2014）、Ampere（2020）、Ada Lovelace（2022）等。

在同一代架构有着统一的微架构设计（如内存控制器、SM 单元结构），支持特定特性（如 Maxwell 引入动态显存压缩，Ampere 集成第三代 Tensor Core）。

第二个是核心名。

核心名代表的是架构下的具体实现。如GA102（Ampere 架构下的高端核心）、AD104（Ada 架构下的主流核心）。

一般来说，同一架构可衍生多个核心，差异在于SM数量，CUDA 核心数量、显存带宽等参数。例如：

GA102（RTX 3090/3080）：10752 CUDA 核心，384-bit 显存位宽。
GA104（RTX 3060 Ti）：4864 CUDA 核心，256-bit 显存位宽。

第三个是大家最熟悉的产品名

产品名面向市场的具体型号，也是大家最熟悉的 GPU 的概念，如GeForce 8800 Ultra、RTX 4090、GTX 1660 Super。

同一核心可衍生多个产品，通过频率、显存容量等差异化定位。例如 GA104 核心下包括 RTX 3070、RTX 3060 Ti 等多个不同的产品。

RTX 3070：8GB GDDR6，2912 MHz 频率。
RTX 3060 Ti：8GB GDDR6，1665 MHz 频率。

给大家推荐一个很好的网站 https://www.techpowerup.com/。在这个网站里可以根据架构名、核心名、产品名来检索英伟达各个 GPU 的详细数据。飞哥的分享中很多数据都来源于这个网站。

2.1 Tesla G80 架构

Tesla 架构包括 G80、GT200 等核心。GT200 是 G80 的改进款，CUDA 核心数、显存容量都有进一步的提升。产品型号中 GeForce 8800GTX 是基于 G80 核心。Tesla C1060基于 GT200 核心。

其中 Tesla G80 作为第一代通用计算产品，设计了 128 个 CUDA 核心数。以下是我找到的 G80 GPU 芯片的外观。

图片

我们来看下 Tesla G80 架构框图。

架构框图是一种用于展示 GPU 内部硬件组成结构、各组件之间的连接关系以及数据流向的示意图，有助于大家理解 GPU 的工作原理和内部架构。

图片

在上图中我们可以看到，Tesla 架构主要的核心部件是包含了8 个 TPC 和 8 个 ROP 光栅操作单元。

1）TPC texture/process cluster 纹理 / 处理簇

TPC 纹理 / 处理簇首次实现了将传统 GPU 中分离的纹理处理单元和计算单元整合到同一模块中。这就是前面我们提到的（Unified Shader Architecture）的含义。

通过 TPC 内各模块的紧密协作，首次实现了图形渲染与通用计算的硬件资源共享。既满足图形渲染的实时性需求，又支持 CUDA 编程的并行计算。这也就是统一着色架构（Unified Shader Architecture）的核心。

2）ROP Render Output Pipeline 渲染输出流水线

ROP 渲染输出流水线比较简单，它负责将 TPC 处理后的数据写入帧缓冲区。

2.2 每个TPC 内部结构

在 Tesla 架构中，TPC 是最核心的组件。让我们放大来看下它内部的结构。

图片

我们看到每个 TPC 内包含了一个 Gepmetry controller、一个SMC、两个 SM 和一个 Texture unit。

1）Geometry Controller 几何控制器

在 Tesla 架构之前，顶点处理由专用的顶点着色器单元完成。统一架构后，Geometry Controller 成为 TPC 的一部分，负责顶点处理。通过共享指令缓存和寄存器文件，提升了硬件利用率。

2）SMC 流多处理器控制器

Tesla 采用SIMT（单指令多线程）架构，以 32 线程为一个 Warp 进行调度。

SMC 全称是 Streaming Multiprocessor Controller 流多处理器控制器，负责线程调度与指令分发。SMC 管理 TPC 内两个 SM 的执行队列，将 CUDA 程序中的线程块（Thread Block）分配给 SM 执行，并确保指令以 Warp（32 线程组）为单位高效执行。

3）SM 流多处理器

SM 全称是 streaming multiprocessor 流多处理器。是并行计算的核心。由于 SM 比较重要，我们单独再开一个小节来介绍它。

2.3 SM 中的缓存单元

每个 SM 都包含有 I cache、C cache 、共享内存。

图片

1）I cache 指令缓存

G80 中包含16KB 的 I cache 指令缓存，用于存储从全局内存加载的 CUDA 核函数指令，减少重复读取显存的开销。指令缓存命中率高时（通常 > 95%），可将指令获取延迟从数百周期降至个位数周期。

2）C cache 常量缓存

G80 中包含 8 KB 的 C cache 常量缓存，用于加速常量数据的访问。

在 CUDA 编程中通过 _constant_ 修饰符声明的数据，例如物理常数、变换矩阵等在核函数执行前会从主机端复制到 GPU 常量 Ccache内。例如下面代码中的 coeff 就是利用了这个特性。

__constant__ float coeff[1024];  // 声明常量内存

__global__ void kernel(float* data) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    data[idx] *= coeff[idx % 1024];  // 从常量缓存读取数据
}

这样 GPU 核在运算的过程中访问 coeff 常量数据时，可直接从 C cache 获取，避免绕行全局显存，性能也大大提升。

3）共享内存

还包含一个 16 KB 的 share memory 共享内存，用于线程块内的线程间通信和数据共享。在 CUDA 编程中，通过 __shared__ 关键字声明使用共享内存。

__global__ void kernelFuncA(float *A, float *B, float *C, int N) {
    // 声明共享内存数组
    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE]
    
    ....
}

2.4 SM 中的运算单元

每个 SM 还包含 8 个流处理器（SP，即 CUDA 核心）、两个 SFU 。

图片

1）SP 流处理器

每个 SM 内包含了 8 个 SP streaming processor。这也就是我们常说的 GPU 的核。所以，Tesla 架构中包含 8 TPC × 2 SM × 8 SP = 128 个核。后来英伟达在 2010 年发布Fermi 架构时开始使用 “CUDA Core” 来代指 SP。

每个 SP（单精度处理单元）每时钟周期最多可完成两次单精度 MAD（乘加）运算。

MAD（Multiply-Add）是一种常用的计算操作，指将两个数相乘后再与第三个数相加，即 “a×b + c”，在图形渲染、科学计算等领域应用广泛。但是硬件的 MAD 操作上是分两步执行：先乘法，再加法。中间结果会被舍入并存储，可能引入舍入误差。而且随着计算的次数越多，误差累积到后面会越来越大。

后来的 GPU 硬件将 MAD 改进成了 FMA（融合乘加）。将乘法和加法融合为一个操作，只进行一次舍入。在计算 “a×b + c” 操作时一步完成，在处理浮点数时结果更精确。

2）SFU 特殊函数单元

SFU 全称是 Special Function Unit 特殊函数单元，是流多处理器（SM）的核心组件之一。专门用于执行超越函数（Transcendental Functions）等较复杂的数学运算。例如：

三角函数：如 sin()、cos()、tan()。
指数与对数：如 exp()、log()、pow()。
倒数与平方根：如 1/x、sqrt(x)。

根据官方资料显示，每个 SFU（特殊功能单元）每时钟周期最多可完成 4 次 SF 操作。这些单元每时钟周期还能处理4 次单精度浮点乘法。

Tesla 架构中的 SFU 通过硬件专用化，显著加速了科学计算和图形渲染。值得一提的是后来 2017 年的 Volta 架构中新增 Tensor Core 也分担了不少复杂数学运算工作。

3）MT issue 多线程指令发射

MT issue 指的是多线程指令发射（Multi-Threaded Instruction Issue）。其核心思想是通过多线程隐藏延迟。

MT issue 实时监控各执行单元（ALU、SFU等）的状态，当一个线程束（Warp）因内存访问延迟或其他原因暂停时，SM 立即切换到另一个就绪的 Warp 执行，避免执行单元空闲。

好了，以上就是 Tesla 架构的基本内容。

三、GeForce 8800 Ultra 算力

在衡量一块显卡的性能时，我们经常会采用峰值FLOPS（Floating-point Operations Per Second）指标。该指标表示每秒浮点运算次数。它是评估高性能计算（HPC）、人工智能、图形渲染等需要大量浮点计算场景的核心参数。

那么对于首代采用 Tesla G80 架构核心的 NVIDIA GeForce 8800 Ultra 这一款产品，它的 FLOPS 算力如何呢？让我们来进行一个简单的计算。

先来看单个 SM 每个时钟周期的 FP32 运算次数

总共有 8 个 SPs，每个 SP 每个周期可执行一次 FP32 MAD 操作（包含两次运算）
总共有 2 个 SFUs，每个 FPU 每个周期也可执行 4 次 FP32 运算

那每个 SM 每个时钟周期可执行的 FP32 运算次数 = 8 SPs * 2 (MAD) + 4 * 2 SFUs = 20 次

每块 GPU 的计算公式如下：

GPU 算力 = Shader工作频率 × SM 数量 × 每个 SM 每个时钟周期的 FP32 运算次数

从 techpowerup 站点上我们找到了 G8800 Shader Clock 是 1512 MHz 。数据地址：https://www.techpowerup.com/gpu-specs/geforce-8800-ultra.c195。则可以算得 GeForce 8800 Ultra 的 FP32 的算力是 387.1 GFLOPS。

GPU 算力 = Shader工作频率 × SM 数量 × 每个 SM 每个时钟周期的 FP32 运算次数
         = 1.512 GHz * 128 * 20
         = 387.1 GFLOPS

这个算力虽然和今天动辄几十 TFLOPS（等于几万 GFLOPS）等算力没法比。

但是和同时期的 CPU 相比较起来，提升幅度非常之大。CPU 侧重单线程效率和复杂指令处理，核心数少且浮点单元占比低。例如，Core 2 Duo 的浮点单元仅占芯片面积的 10%。

我们以 2007 年的 Intel Core 2 Duo E670为例来作为对比。该 CPU 是双核，主频是 2.66 GHz，每个核心每周期可执行 2 次单精度浮点运算（SSE2 指令）。理论 FP32 算力 = 2核心 × 2 FLOPS/core × 2.66e9 Hz = 10.64 GFLOPS。这还没算因为内存带宽限制影响因素。

GeForce 8800 Ultra 的 FP32 峰值算力大约是同时期消费级 CPU 的 30 倍之多。

四、CUDA 计算统一设备架构

在 2006 年除了 Tesla 架构统一着色架构只是硬件基础。接着 NVIDIA 在 2007 年推出的 CUDA 1.0，全称是 Compute Unified Device Architecture 计算统一设备架构。首次实现 GPU 的高效通用计算能力，打破了 GPU 只能处理图形渲染的限制。

4.1 简单的 CUDA 例子

为了方便理解，我们直接从一个简单的例子来看。

以下是一个使用 CUDA 1.0 语法编写的简单向量加法示例。这个示例展示了 CUDA 1.0 的基本编程模型，包括核函数定义、内存管理和线程索引计算。

#include <stdio.h>
#include <cuda_runtime.h>

// CUDA 核函数定义：向量加法
__global__ void vectorAdd(const float *a, const float *b, float *c, int n) {
    // 计算全局线程索引
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < N) c[i] = a[i] + b[i];
}

int main() {
    // 分配主机内存并初始化
    float *h_a, *h_b, *h_c;
    h_a = (float*)malloc(size);
    h_b = (float*)malloc(size);
    h_c = (float*)malloc(size);
    ...
    
    // 分配设备内存
    float *d_a, *d_b, *d_c;
    cudaMalloc((void**)&d_a, size);
    cudaMalloc((void**)&d_b, size);
    cudaMalloc((void**)&d_c, size);
    
    // 将数据从主机复制到设备
    cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);
    
    // 设置线程块和网格维度
    // CUDA 1.0 要求每个线程块的线程数不超过512
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    
    // 执行核函数
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);
    
    // 将结果从设备复制到主机
    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);
}

在上面的代码中，大概包含了如下三个关键步骤完成 CUDA 程序执行。

第一步：GPU内存申请和复制

先调用 cudaMalloc 完成 GPU 全局内存申请。接着再调用 cudaMemcpy 函数将数据从主机复制到 GPU 设备。

第二步：设置线程快和网格维度，并让 GPU 执行核函数

在 CUDA 中引入了核函数（Kernel）的概念。所谓核函数，是用 __global__ 修饰的函数，由 CPU 调用并在 GPU 上并行执行。

__global__ void vectorAdd(const float *a, const float *b, float *c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < N) c[i] = a[i] + b[i];
}

定义完后，在 C 语言中就可以调用它，让 GPU 执行这个函数。

vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);

第三步：将结果拷贝回 CPU 内存中。

// 将结果从设备复制到主机
    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);

4.2 CUDA 1.0 原理

CUDA 核心解决的问题就是提供了 C 语言编程接口，可以让用户像使用 CPU 一样方便地使用 GPU。

CUDA 提供主机 - 设备（Host-Device）编程范式。CPU 作为主机（Host）负责控制逻辑，GPU 作为设备（Device）执行并行计算。设计实现了三大块功能

GPU 内存的管理
核函数定义
GPU核的抽象

我们来分别看下这三块功能：

1）GPU 内存的管理。

平时我们在 C 语言中使用 malloc 等方式申请的内存是位于和 CPU 相连接的 DDRx 内存中的。那么要进行 GPU 编程的第一步，就是需要有能力操作 GPU 中的 GDDRx 显存（新的英伟达的 GPU 已经不再使用 GDDRx，而是开始使用 HBM 颗粒，这个我们将来再将）。

图片

CUDA 中定义的 GPU 内存管理相关函数有：

cudaMalloc：申请 GPU 全局内存
cudaMemcpy：在 CPU 与 GPU 之间互相拷贝的方式传输数据
cudaFree：释放 GPU 全局内存

在 GPU 内存中准备好了数据之后，下一步就可以让 GPU 来调度自己的 SP 核来处理数据了。

2）核函数定义。

在 CUDA 中引入了核函数（Kernel）的概念。所谓核函数，是用 __global__ 修饰的函数，由 CPU 调用并在 GPU 上并行执行。

__global__ void vectorAdd(const float *a, const float *b, float *c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < N) c[i] = a[i] + b[i];
}

3）GPU 核的抽象。

为了方便用户使用，CUDA 对硬件进行了抽象。将 GPU 的流处理器抽象为线程（Thread）、线程块（Thread Block）和线程束（Warp），开发者无需关心图形 API 即可调用 GPU 算力。其中

线程（Thread）：最小执行单元，执行相同指令（SIMT）。
线程块（Thread Block）：一组协作线程，共享内存并可同步。
网格（Grid）：多个线程块的集合，可分布在多个 GPU 流多处理器（SM）上。

在上面的demo代码的 vectorAdd 核函数的实现中。

__global__ void vectorAdd(const float *a, const float *b, float *c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;  // 计算全局线程索引
    if (idx < n) {c[idx] = a[idx] + b[idx];}
}

threadIdx.x 表示线程在块内的局部索引（x 维度）。每个线程：执行相同的向量加法指令 c[idx] = a[idx] + b[idx]，但处理不同的数据（通过 idx 区分）。

另外threadsPerBlock 定义每个块的线程数量是 256 个。

int threadsPerBlock = 256;  // 每个线程块包含256个线程

每个线程块（Block）被划分为多个连续的 Warp。每个 Warp 固定大小为32。若线程块大小为 256，则被划分为 8 个 Warp（256 ÷ 32 = 8）。每个 Warp 由一个 SM 来处理。SM 内的调度器和执行单元以 Warp 为粒度操作调度。

Tesla GPU 采用 SIMT 架构。每个 SM 管理 24 个 Warp 的资源池。每个 Warp 支持 32 个线程。所以 SM 在零调度开销的情况下可以最大支持 768 个线程并行。

blocksPerGrid 根据数据总量 n 和每个块的线程数 threadsPerBlock 动态计算所需线程块数。

int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;  // 计算所需块数
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);  // 启动核函数

网格中的不同线程块可被分配到 GPU 的不同流多处理器 SM 上执行。我们在 3.2 中看到 Tesla 架构中有 8 个 TPC，每个 TPC 下有 2 个 SM，总共有 16 个 SM。也就是说，我们这段简单的代码，最多是由可能由 16 个流多处理器 SM 同时来处理的。

再结合第二节我们学到的 Tesla 的内部结构，我们可以看到，CUDA 通过线程、线程块、网络分别对应了 GPU 硬件中的流处理器 SP、流多处理器 SM以及整个 GPU。使得用户可以方便地通过这些概念来操作 GPU 硬件。

CUDA软件抽象	对应硬件组件	Demo 中的参数
线程（Thread）	流处理器（SP）	threadIdx.x
线程块（Block）	流多处理器（SM）	threadsPerBlock = 256
网格（Grid）	整个 GPU	blocksPerGrid = ceil(n/256)

当然了早期的 1.0 功能还比较薄弱，例如仅支持 FP32，需手动管理内存和线程层次，学习门槛高。但仍然不妨碍它成为一个划时代的技术产品。

CUDA 后面的版本里进行了很多完善，当然这些完善是包括硬件架构和软件的统一升级。例如：

CUDA 2.0（2008）：引入动态并行、原子操作。
CUDA 5.0（2012）：支持双精度浮点（FP64），优化 Kepler 架构。
CUDA 11（2020）：引入安培架构，支持张量核心和混合精度计算。

总结

虽然说英伟达的股价是在 2020 年之后开始逐步起飞的，但真正的历史原因却要追溯到 2006 年 Tesla 架构的发布和 CUDA Toolkit这两大技术突破中。

之所以说 Tesla 是英伟达的第一代通用计算 GPU，主要是以下两个原因：

第一、通过统一着色架构（Unified Shader Architecture）将传统分离的顶点着色器、像素着色器、几何着色器统一为可编程的流处理器（Streaming Processor）。最重要的是该流处理器也可执行与图形无关的计算指令，为通用计算提供了基础。

第二、接着 NVIDIA 在 2007 年推出的 CUDA 1.0，全称是 Compute Unified Device Architecture 计算统一设备架构。CUDA 提供主机 - 设备（Host-Device）编程范式。CPU 作为主机（Host）负责控制逻辑，GPU 作为设备（Device）执行并行计算。设计实现了GPU 内存的管理、核函数定义、GPU核的抽象三大功能。

通过 CUDA 支持使用 C 语言进行 GPU 编程，改变了以往 GPU 只能用于图形渲染的局限，使开发者能够利用 GPU 的并行计算能力处理通用数据，开启了 GPU 通用计算的时代。

此时，其它家竞争对手尚未大规模投入 GPGPU （General-Purpose computing on Graphics Processing Units ，指的是用 GPU 执行通用计算任务）领域，而英伟达却通过 Tesla 架构和 CUDA 提前卡位。CUDA 开发工具链、优化库（如 cuBLAS、cuDNN）和开发者社区支持，的推出吸引了大量开发者，开始建立起围绕 GPU 计算的开发生态。到 2010 年，CUDA 已成为 GPU 计算的事实标准。

就这样，英伟达通过 Tesla 硬件架构和 CUDA。首次实现 GPU 的高效通用计算能力，打破了 GPU 只能处理图形渲染的限制。实现了从游戏显卡到通用计算的华丽转身。

这一决策彻底改变了英伟达的命运，使其从一家普通的 “显卡公司” 转型为 “计算平台巨头”，并推动了今天人工智的能的革命。