目录

1960s~1999：GPU 的诞生：光栅化（Rasterization）3D 渲染算法的硬件化实现之路

学术界算法研究历程

1960s：光栅化技术理论研究的核心思想起源于计算机图形学的初步探索，研究者的目标是解决几何投影与像素覆盖的问题，即：如何将 3D 几何体（如正方体）投影到 2D 屏幕空间上。

1970s：**扫描线算法（Scanline Algorithm）**被提出，通过逐行填充多边形边界内的像素，显著提升了光栅化效率。同时，几何投影理论（如透视投影与正交投影）逐渐形成，奠定了光栅化流程的数学基础。

1980s：随着图形化桌面操作系统的研究发展，计算机图形学进入实用化阶段。此时光栅化技术和光线追踪（Ray Tracing）技术形成了路线竞争，并最终胜出。因为光线追踪技术的实时渲染需要依靠大量算力，当时硬件条件并不具备，直到 2018 年 Turing GPU 才将光线追踪技术带进了市场。

1980s 后期：**深度缓冲（Z-Buffer）**技术概念的雏形出现，可应用于解决多个几何物体前后遮挡的问题，成为光栅化的标准配置。

1990s 前期：纹理映射（Texture Mapping）技术成熟，光栅化算法开始支持复杂表面细节的渲染，同时抗锯齿（Anti-Aliasing）技术也通过多重采样的方式优化了边缘锯齿的问题。

2000s：可编程着色器（Shader）技术成熟，使光栅化管线（Pipeline）从固化的 GPU 硬件功能走向灵活可编程的趋势，支持逐像素光照计算（如 Phong 着色、法线贴图），这极大提升了渲染效果。

工业界产品研发历程

1981 年，现代显卡的雏形：IBM 推出的 5150 PC 搭载了 MDA（Monochrome Display Adapter，单色显示适配器）和 CGA（Color Graphics Adapter，彩色图形适配器）这 2 款显示适配器，奠定了计算机图形显示的基础。

1989 年，第一代 VGA Card（显卡）设备：随着 Microsoft Windows 图形化操作系统的发展而诞生的 VGA Card 只提供 2D 图像显示功能，支持在屏幕上输出 640×480 分辨率、256 彩色的图像。另外，除了显示之外的图形运算依旧由 CPU 完成。

1991 年，第二代 Graphics Card（图形加速卡）：专用于支持 Windows 图形界面。通过一颗专用芯片来同时处理 2D 显像和图形运算，卸载了 CPU 的图形运算。让更多的 CPU 资源用于运行 Windows 操作系统，时期运行非常流畅，也加速了 Windows 的普及，让 PC 走进了图形化界面时代。

1994 年，第三代 Video Card（视频加速卡）：不仅仅是 2D 图形显像，更一步支持了针对视频显像数据流的编解码加速。

1996 年，第四代 3D Accelerator Card（3D 图形渲染加速卡）：这是 NVIDIA 推出的一款划时代的产品，内含 TNT/TNT2 芯片的 3D 显卡，用硬件的方式实现了光栅化算法（Rasterization Stage），但依旧依赖 CPU 来完成顶点处理（Geometry Stage）的运算。

1999 年，第五代 GPU（图形处理器）：GeForce 256 是 NVIDIA 推出的另一款划时代的产品，第一次提出 GPU 芯片的概念。面向发展迅猛的 3D 游戏图像处理加速市场，在游戏领域，第一次让 GPU 上升到了与 CPU 同等的高度，3D 渲染显像技术走向普及。

NVIDIA GeForce 256 作为世界上第一块 GPU，距今（2025）已有 25 年历史。

光栅化技术原理

光栅化技术的第一性原理是 “投影”，即：选取一个视角（Camera，摄影机），然后将 3D 几何物体透视投影到画布上，最后将投影像素化的过程。更进一步的，除了虚拟摄像头，还需要具有虚拟光源，并且当多个物体投射到同一张画布上时，还需要通过一系列的采样和插值算法来确定画布上像素点的取值。

光栅化技术的软件实现：OpenGL 3D 渲染管线设计

3D 渲染管线（Graphic Pipeline）能够实时的将 3D 图形渲染为 2D 屏幕中进行显示，并且参考 CPU 流水线，GPU 也通过流水线的方式，大大加快了实时渲染的速度。但即便如此，渲染一个复杂的游戏或电影场景依旧需要非常庞大的算力和时间。

粗略的，可以将 3D 渲染管线分为 “几何处理” 和 “光栅处理” 这 2 大阶段。

1. 几何处理：把 3D 坐标转换为 2D 坐标。
2. 光栅处理：把 2D 坐标转变为实际有颜色的像素方块。

更详细的，这 2 个阶段又可以细分为多个工作模块。这些工作模块在 GPU 中有相应的硬件模组进行处理，并且针对这些工作 GPU 也提供了不同的可配置性或可编程性。即：有些工作的硬件功能固化的，而有些工作的硬件功能是可编程的，这也是 CUDA 诞生的基础。例如：

1. 顶点数据的输入与处理（顶点着色器，可编程）：完成坐标变换与属性传递。
2. 曲面细分（曲面细分着色器，可编程）：完成动态曲面细分。
3. 图元组装（几何着色器，可编程）：完成重新生成图像。
4. 裁剪剔除（不可编程）：剔除不可见部分。
5. 屏幕映射（不可编程）：NDC 到屏幕像素。
6. 图元光栅化（三角形设置与遍历，不可编程）：生成片元。
7. 像素着色（片元着色器，可编程）：计算颜色与特效。
8. 片元筛选（逐片元操作）：深度/模板测试与混合。
9. 屏幕图像输出（帧缓存）：存储并显示最终的屏幕 2D 图像。

值得注意的是，根据 GPU 型号和软件框架（如 OpenGL）的不同，Pipeline 的设计并不完全一致。如下图是早期 OpenGL 的一个 Pipeline 软件设计，以及早期 RealityEngine 显卡芯片的硬件设计。

1. 顶点处理（Vertex Processing）

OpenGL Application 通过 CPU 调用 DrawCall 函数，启动一个 3D 渲染任务，向 GPU 输入一系列 3D 顶点坐标及其属性信息，包括：顶点坐标、纹理坐标、顶点法线和顶点颜色等顶点属性。为了让 OpenGL 明白顶点数据构成的是什么图元（Primitive），在 Application 中需要使用绘制指令传递相对应的图元信息，包括：点（GL_POINTS）、线（GL_LINES）、线条（GL_LINE_STRIP）、三角面（GL_TRIANGLES）等参数。

GPU 的顶点着色器（Vertex Shader）会对每个顶点坐标进行一系列的变换，包括：几何变换、视图变换、投影变换等，将原始 3D 坐标从 Local Space 映射到 Global Space，然后根据视角的不同，再从 Global Space 转换为 Screen Space，最后再从 3D 投射到 2D 坐标。

这之间主要计算方式就是线性代数中的矩阵乘法，并且考虑到齐次坐标（homogeneous coordinates），这通常是一个四维矩阵乘法运算。最终，顶点着色器将这些顶点坐标全部转换世界坐标、观察坐标和裁剪坐标等等。

2. 三角处理（Triangle Processing）

三角处理，也称为图元处理（Primitive Processing），就是将前面得到的一系列 2D 坐标点组合在一起构成许多三角形，继而由多个三角形之间的拼接构成一个新的 3D 轮廓。除了最基本的三角处理，现在 GPU 和 OpenGL 还支持以下新特性：

1. 曲面细分着色器（Tessellation Shader）：利用镶嵌化处理技术动态地将三角形进一步细分，增加物体表面的三角面的数量，使曲面更圆滑。曲面细分由外壳着色器（Hull Shader）、镶嵌器（Tessellator）和域着色器（Domain Shader）构成，其中外壳着色器和域着色器是可编程的，而镶嵌器是有硬件固化的。
2. 几何着色器（Geometry Shader）：用于将输入的图元（点或线）扩展成几何多边形。
3. 图元组装与裁剪剔除：将输入的顶点组装成指定的图元。同时将超出画布或位于背面的图元剔除或裁剪成新的三角形，以减少进入光栅化的图元的数量，加速渲染过程。
4. 屏幕映射：透视除法和视口变换。

3. 光栅化（Rasterization）

经过图元组装以及屏幕映射阶段后，GPU 已经将物体的 3D 坐标转换为了 2D 的屏幕坐标。而光栅化，也称为栅格化，本质是一个离散化运算，就是将前面得到的每个三角形离散为一个个像素方格，因为屏幕所显示的图像就是由一个个像素点组成的。

在离散过程中，光栅化会利用图元的信息和一系列插值算法来确定图元像素所覆盖的片段（Fragments），并使用顶点属性的插值获取到这些片段的属性信息。这些信息用于后续的片段处理。

4. 片段处理（Fragment Processing）

片段处理，又称为纹理帖图（Texture Mapping）。由片段着色器（Fragments Shader）根据光栅化后的顶点属性信息（如着色设计）为每个像素点进行上色。需要注意是，片段是像素的候选者，只有通过后续的测试，片段才会成为最终显示的像素点。

另外，根据 GPU 厂商不同，片段着色器（Fragments Shader）也可以称为像素着色器（Pixel Shader）或 TMU（Texture Mapping Unit）。

5. 帧缓存（Frame Buffer）输出

当 1 帧图像的所有并行计算得出的 Fragments 送到显存帧缓冲区进行合并，合并的过程中会使用到 Z-buffer 等操作，最终构成像素级的图片，输出到显示器上。

光栅化技术的硬件实现：GPU 芯片架构设计

1996 年，NVIDIA 虽然已经通过 TNT 芯片实现了 3D 渲染管线中的光栅处理阶段（Rasterization Stage），但依旧依赖 CPU 来完成几何处理的（Geometry Stage）的运算。如下图所示，在 TNT 芯片中集成了 Rasterizer、Texture、Raster Operations Unit，Frame Buffer 等模块。

1. CPU 负责所有的几何处理，包括：Vertex Shader 和 Geometry Shader，而将生成的 Primitive 2D Triangles 和相应的纹理材质加载到 Main Memory 中。
2. 2D Triangles 被 TriangleCMD 通过 AGP bus 发送到 TNT 芯片，通过 Rasterizer 转换成扫描线请求并发送到 TMU（Texture Mapping Unit）。
3. TMU 将 fragment 填充好之后，通过 Raster Operations Unit 输出到 Frame Buffer。
4. Frame Buffer 利用 Z-Buffer test 技术在 GPU Memory 中保存 RGBA 和 z 值，并在完成了混合处理等特效功能后，转换成 VGA 信号给显示器。
   

3 年后的 1999 年，NVIDIA 发布 GeForce 256 显卡，并第一次提出 GPU（Graphics Processing Unit，图形处理器）芯片的概念。之所以命名为 GPU 是因为这块芯片已经硬件实现了 3D 渲染管线的整个流程，而不再需要 CPU 的参与。

如下图所示，G256 在 TNT 的基础上还集成了 T&L（Transform & Lighting Unit，光照变化单元）模块，这是一个固化的 32bits 浮点矢量计算单元，实现了顶点处理和几何变换等功能，而 CPU 只需要负责将 3D Triangles（描述 3D 空间的三角形）。

2000~2006：CUDA 的诞生：从 GPU 专用计算走向 GPGPU 通用计算

硬件芯片设计基本原则

硬件芯片有几十种大类，上千种小类，例如：处理器芯片、存储器芯片、逻辑器芯片等等。而芯片的使命就是要运行软件程序，例如：

• 中央处理器（CPU）：运行 Windows、Linux 等复杂操作系统；
• 图形处理器（GPU）：不会加载操作系统而是直接运行 OpenGL、HPC、AL 等框架程序；
• 神经网络处理器（NPU）：直接运行 DL（深度学习）相关的程序；
• 微控制处理器（MCU）：运行实时操作系统或者直接运行某个特定程序。

处理器芯片设计过程复杂，包括：程序特征分类与提取、微架构设计空间探索、PPA 多目标优化等多个环节。其中最关键的就是要从需求出发，从处理器所承载的程序特征出发，决定了芯片设计的方向。

而处理器微架构设计优化的一个重要思路就是发现程序行为中的共性特征并进行加速，包括：缓存、流水线、并行、预取等，都对应着程序行为的共性特征，比如缓存和预取就需要充分利用程序的局部性特征。可以再进一步细化为时间局部性与空间局部性，对应到程序的顺序执行、条件判断和循环。

以分支预测单元 BPU（Branch Prediction Unit）为例，其负责根据分支指令执行的历史记录，去预测后续分支的走向，从而提前预取、执行对应方向上的指令。可见，BPU 预测的准确率直接影响了处理器的性能：

• 当 BPU 预测准确率高，则处理器流水线的空泡（Stall）就比较少，甚至完全消除。
• 当 BPU 出现预测错误，则不仅已执行的错误路径上的指令都被浪费，而且还需要冲刷流水线等来保证后续执行的正确性，这降低了处理器性能。

而 BPU 准确率高低的关键就在于发现程序执行过程中的分支行为特征。

业务需求驱动硬件芯片可编程

NVIDIA GeForce 256 加入了 T&L 模块之后使其成为了世界上第一张完全硬件实现了完整 OpenGL 3D 渲染管线的 GPU 设备。

而后随着需求侧 3D 游戏的迅猛发展，游戏开发者们需要实现更复杂的光照、阴影、纹理映射等效果。完全硬件固化的渲染管线逐渐无法满足越来越丰富的 API 需求和越来越复杂的渲染管线设计。所以，供给侧在 2001~2006 这几年间的主题之一就是提供更好的可编程 GPU 设备。

• 软件层 DSL（Domain Specific Language）：着色器编程语言和 API 标准化，包括：OpenGL（GLSL）、DirectX（HLSL）、NVIDIA（C for graph）等。允许开发者编写顶点着色器（Vertex Shader）、片段着色器（Fragment Shader）和几何着色器（Geometry Shader）等程序，以控制图形的渲染过程（如顶点变换、光照计算、像素着色等）。
• 硬件层 DSA（Domain Specific Architecture）：形成了 NVIDIA 和 ATI 双寡头竞争市场，逐渐实现了 VS、GS、PS 等芯片器件的可编程性。
  

Vertex Shader 可编程

2001 年，NVIDIA GeForce 3X 是第一款支持可编程 VS 的 GPU。如下图所示，G256 的 T&L Unit 和 Texture Unit 被替换成为 Vertex Shader 和 Texture Shader 模块。

• Vertex Shader：是一个可编程的 32bits 浮点矢量计算单元（User-Programmable Vertex Engine），通过一系列的 SIMD（Single Instruction Multiple Data，单指令多数据流）浮点计算指令，向开发者暴露了一部分顶点计算的可编程能力。但仍不支持分支等控制单元功能（no flow control）。
• Texture Shader：仅支持配置操作，仍不支持编程。

值得注意的是，从光栅化算法的基础原理上就决定了 3D 图形的每一个顶点的处理之间并没有依赖，所以可以采用并行计算的方式。这也决定了 GPU 芯片设计之初就是为了并行计算而生的。

在 SIMD 架构的处理器中，一个控制单元可以同时向多个运算单元发送相同的指令，并且每个运算单元支持操作不同的数据集。例如，一条加法指令可以同时对多个数据集进行加法运算。相对于 MIMD，SIMD 更适用于并行处理数据密集型任务，例如：图像处理、视频解码、HPC 和 AI 等。

Pixel Shader 可编程

2003 年，NVIDIA 发布 GeForce FX，如下图所示，继续增加了 Pixel Shader 模块，这也是一个可编程的 32bits 浮点矢量计算单元。此外，VS 模块也增加了分支控制能力，支持单个程序最大指令数 65536 条，循环深度 256 个，还支持了 static and dynamic flow control。

2004 年，伴随着 NVIDIA GeForce 6X 的发布，VS 和 PS 都实现了完全的可编程性。

G6X 的 VS 和 PS 模块的架构如下图所示：

• 控制单元：都支持了完整的分支、循环、预测等控制能力。
• 运算单元：FP32 ALU（单精度浮点数），满足顶点变换、光照模型和纹理映射等运算任务对数值精度的要求。
• 存储单元：L1、L2 cache
• 功能单元：TMU 等等
  

Geometry Shader 可编程

2006 年，NVIDIA 再把 GS 可编程模块补全。从 1999~2006 年的 7 年间，NVIDIA 终于把 VS、PS、GS 可编程模块集齐，作为第一个统一支持了多种 3D 渲染 DSL（Cg、GLSL、HLSL）的 DSA 硬件平台，并同时提出了 CUDA 的前身 USPA（Unified Streaming Processing Array）—— Can be unified hardware。

性能需求驱动硬件芯片统一架构

在 NVIDIA 可编程化了 VS、PS、GS 芯片模块之后发布的 GeForce 6X 到 GeForce 7X 的过程中，如何平衡三者之间的成为了首要问题。

如下图所示，在一个 non-unified architecture（非统一架构）的 GPU 中各个芯片模块的算力使用率并不均衡。例如：在曲面较多的复杂场景中，有大量的曲面细化处理，会生成大量的三角形需要被 VS 处理，那么 VS 就会跑满，而后续环节的 PS 却空闲。

因此如何将 VS、PS、GS 统一为一个 Unified Shader 就成为了工业界的主题之二。

GPGPU = Tesla Architecture + CUDA Platform

2001 年，斯坦福大学 Bill Dally 教授团队在《IEEE Micro》（处理器微结构旗舰期刊）上发表了一篇题为《Imagine:Media Processing with Streams》的论文正式介绍一种可加速多媒体应用的流处理器（Stream Processor）架构。这项工作很快引起了 NVIDIA 的关注。

2003 年，Dally 教授担任 NVIDIA 的顾问，参与 Tesla Architecture GPU 的微架构设计，研究在 GPU 中加入 SP。与此同时，斯坦福大学的博士生 Ian Buck 在导师 Pat Hanrahan 教授（2019 年图灵奖得主）的指导下开展如何方便且高效发挥 GPU 算力的研究。

2004 年，Buck 发表了一篇题为《Brook for GPUs:Stream Computing on Graphics Hardware》的论文，通过 Extensive C lang 为 GPU 设计了一套支持流编程（Stream Programming）语言，并实现了一个 BRCC 编译器，运行时系统 Brook，利用当时的 NIVDIA GeForce 6X GPU 实现了 SGEMV（矩阵运算）、FFT 等 5 种通用计算的算法，从而能让开发者在 GPU 上也像在通用 CPU 上那样进行编程。于是 Buck 博士毕业后便立刻加入 NVIDIA，带领两位工程师创立了 CUDA 项目。

2006 年 11 月，采用 Tesla Architecture 的第一款 GeForce 8800 GTX GPU 发布，包含 128 个 SP，单精度浮点运算性能达到 345.6GFLOPS，访存带宽 86.4GB/s，远高于同期 CPU 性能。

2007 年，CUDA 1.0 正式发布，全面适配 GeForce 8800（G80）系列 GPU。

随后，UIUC 胡文美教授团队在 GeForce 8800 GPU 上用 CUDA 实现一些程序，性能比通用 CPU 高 10 倍到 400 倍不等，充分展示了 GPU 的高性能与可编程性。

至此，NVIDIA GPU 的统一计算软硬件生态大厦的地基已基本构成，NVIDIA GPU 开始被称为 GPGPU（General Purpose GPU，通用 GPU）。

Tesla 统一运算架构

Tesla 架构是 NVIDIA 第一个实现了 Unified Shader（统一着色器模型）的芯片架构，合并了 VS 和 PS 模块，具有完整的 Transform（坐标转换）、Lighting（光源处理）、Setup（三角形设置）和 Rendering（渲染引擎） 功能。

Tesla 架构的模块组成如下图所示，大致上可以分为控制单元、计算单元、存储单元和片内总线这 4 大类型。

• Host Interface：用于连接主板 PCIe Bridge
• Input Assembler / Vertex Work Distribution：用于顶点坐标输入和顶点处理计算任务分发。
• Viewport/Clip/Setup/Rester/Z-Cull / Pixel Work Distribution：用于光栅化和像素处理计算任务分发。
• Compute Work Distribution：用于一般通用计算任务分发。
• High-Definition Video Processors：高清视频处理器。
• 3 个 Work Distribution 对接 7 个 TPC（Texture Processor Cluster，纹理处理集群）模块，构成了一个可扩展的处理器阵列。
• 1 个 TPC 包含：
  • 1 个 Geometry controller：TPC/VS 先处理 Vertex，当需要 GS 时，由 Geometry Controller 负责再将 Input Assembler 组装好的 Primitive 再次送到 TPC/PS 处理。
  • 1 个 SMC：负责将计算任务拆分打包成 Warp（32 个 Threads） 并交给某个 SM 执行，同时还负责协调 SM 和 Texture Unit 以及外部资源的获取。
  • 1 个 Texture Unit：通过 SMC 和 ROP 实现外存的读写。
  • 2 个 SM（Stream Multi-processor，流多处理器）
• 1 个 SM 包含：
  • 1 个 I-Cache（指令缓存）
  • 1 个 MT Issue（多线程指令获取）：用于把 Warp 任务拆分成一条条指令分配给 SP 处理。
  • 1 个 C-Cache（常量缓存）
  • 6 个 SP（Streaming Processor，流处理器）
  • 2 个 SFU（Special Function Unit，特殊函数的计算单元）：用于特殊函数的计算，每个 SFU 包含了 4 个浮点乘法器。
  • 1 个 Shared Memory（共享内存）：16KB 大小。
• 1 个共享的 L1 Cache（一级缓存）
• Inter-connection Network（访问存储单元）：片内高速总线，在 SM 和 Main Memory 之间快速传输数据。
• Warp Scheduler（线程调度器）
• 4 个共享的 L2 Cache
• 4 个 HBM 显存：超大带宽的高速显存。
• 1 个 Display Interface：连接显示器输出图像。

后面我们会重点介绍与 GPU 编程关系最密切的 SP、Warp Scheduler 以及显存这 3 个部分。

G80 的基本工作流程

1. GPU host Interface 接收来自 CPU 的计算任务，并从 Main Memory 获取计算任务的代码和数据并实现上下文的交换。
2. Input Assembler 采集几何图元（Geometry primitives），包括：点、线、三角形等，然后获取相关的顶点坐标和属性信息，然后输出给 Vertex Work Distribution 模块。
3. Distribution 模块采用 Round-Robin 的方式将 Vertex Work packets 发送到不同的 TPC 处理。
4. TPC 执行 Vertex Shader 程序，执行完成后写到 onchip buffer 然后通过 viewport、clip、setup、raster/zcull 等模块转换成像素片段（Pixel Fragment）。
5. 随后 Pixel Work Distribution 再进行一次调度给 TPC 进行 Pixel Shader 程序处理。
6. 处理完的 Fragments 由 ROP 处理，并最终输出到 Display。

其中，VS 程序只处理 Vertex，而 PS 程序只能处理 Primitive，所以 TPC 处理完 Vertex 后，还需要 Input Assembler 先将 Primitive 组装好之后，再重新送到 TPC 处理。

Streaming Processor（流处理器）

SP 是 GPU 最核心、最基本的计算单元，指令和运算最终都在 SP 上处理。G80 总计集成了 128 个 SP 标量计算单元。如下图所示，每个 SP 包含了：

• 1 个共享的 Fetch/Decode（取指/译码）部件
• 8 个 ALU（逻辑运算单元）
• 4 组 Execution contexts（执行环境）

在处理器的 ALU 运算单元中，通常存在标量（Scaler）、矢量（Vector）和张量（Tensor）这 3 大类不同的数据类型。它们分别是 3 种不同维度的数据，并在计算中以不同的方式进行处理和优化。

• 标量（Scaler）：零维数据，例如整数或浮点数。标量计算处理器的典型如 CPU，具有复杂的控制能力。负责基本的算术逻辑运算（如加减乘除）和控制流操作（如分支判断）。
• 矢量（Vector）：一维数据，例如数组或列表。矢量计算处理器的典型如 GPU，具有强大的并发能力，能够同时对多个数据元素执行相同的操作（如 SIMD），支持浮点数和整数的并行计算，通常用于图形处理、信号处理等需要高吞吐量的场景。
• 张量（Tensor）：多维数据，例如矩阵是二阶张量，彩色图像可以表示为一个三阶张量（高度、宽度、颜色通道）。张量计算处理器的典型如 TPU，用于高效执行矩阵乘法等张量运算，通过硬件优化（如分块计算、并行处理）加速深度学习中的张量操作。

而 G80 中的 SP 则采用了纯粹的标量处理单元，而没有矢量处理单元，如下图所示，SP 支持 Int、FP32 MAD、FP32 MUL 等整数和浮点数标量运算。

之所以如此设计，是因为标量处理单元对于编译器的实现而言更加容易，标量 SP 的编译器完全支持 C 编程，避免程序员要手工处理矢量寄存器（Vector），为开发上层 CUDA 提供了足够的支撑。真如前文所说，矢量处理单元实际上更适用于并行计算，但 NVIDIA 选择标量处理单元的勇气是 CUDA 成功的关键之一。

Warp Scheduler（线程调度器）

GPU 程序编程的计算任务调度单元依旧是 Threads，称为 CUDA threads。但 CUDA Threads 显然不同于 CPU Threads，但可类比。

• 相似点：GPU 同样存在一个 Scheduler 来完成 CUDA Threads 在 SP 中的调度。
• 差异点：GPU Threads 的数量实在比 CPU Threads 多得多得多。

所以 GPU 需要对 CUDA Threads 进行分组打包以便管理，这就是 Warp（线程束）的理念。

• 1 个 SP 内含有 4 组 Execution contexts，支持并发交叉（Concurrent but interleaved）执行 4 条 Instruction Streams（指令流）；
• 1 组 Execution contexts 又可分为 8 个 Context，每个 Context 中运行 4 个 CUDA threads，共计 32 个 CUDA threads，被打包成一个 Warp。

可见，GPU 的并发粒度是 SP，而实际上每个 SP 中也还含有并发计算。每个 Warp 包含了 32 个并发交叉的 Threads，而 Warp Scheduler 则批量调度这些 Threads 的处理。

另外值得注意的是，由于 G80 采用 SIMT 模式，区别于 SIMD，G80 的每个 SM 都有独立的 Branch（分支）能力。每个 Thread 都有自己独立的指令地址和寄存器状态，使得 Threads 的编程更加灵活。在特定场景中，开发者还可以通过 C 代码描述单个 Thread 的执行。

SM 中的 MT Issue（Multi Thread Issuer）部件用于创建、管理、调度和执行这些 Threads。即：当 Scheduler 将一个 Warp 交给 MT Issue 后，它会将其打散成 Instruction 让 SP 和 SFU 将一条指令累计执行 32 次。

CUDA 统一编程平台

CUDA（Compute Unified Device Architecture）编程框架的设计受到了 OpenMPI 的影响，采用 Grid、Block、Thread 递进层级的抽象概念来组织并行计算任务的定义和调度，与底层 GPU 硬件实现了解耦。

Kernel 概念：GPU 和 CPU 的协同

CUDA Application 的执行区域分为 2 个部分：

1. CPU（Host Code）：CUDA Kernel 函数声明一个并发任务。
2. GPU（Device Code）：实际的 CUDA 计算任务。

每当 CPU 遇到需要并行计算的任务，则将要做的运算组织成 kernel，然后丢给 GPU 去执行。当然任务是通过 CUDA 系统来丢的，CUDA 在把任务正式提交给 GPU 前，会对 kernel 做些处理，让 kernel 符合 GPU 体系架构。一个典型的 GPU 工作流程如下：

1. CUDA Application 调用 CUDA API 启动一个 Workload。
2. CUDA lib 库，通过 UMD（User Mode Driver），提交 Workload 到 KMD（Kernel Mode Driver）。
3. KMD 写 CSR MMIO，把它提交给 GPU 硬件。
4. GPU 开始工作，并在完成后，将结果 DMA 到 Main Memory，然后发中断给 CPU。
5. CPU 找到中断处理程序，由于 KMD 此前向 OS Kernel 注册过的，所以可以调用它。
6. 中断处理程序找到是哪个 Workload 被执行完毕了，最终驱动唤醒相关的 Application 读取结果。

并行计算任务的定义实在 CPU 完成的，而并行计算任务的调度和运行实在 GPU 中完成的。由于 GPU 和 CPU 的 Main Memory 是分离的，所以 GPU 运算时必须先将 CPU 的代码和数据都传输到 GPU 之后，GPU 才能开始执行 kernel 函数。其中的数据流为：

1. 将 Main Memory 的处理代码和数据复制到 GPU Memory 中。
2. CPU 指令驱动 GPU。
3. GPU 中的每个运算单元并行处理。此步会从 GPU Memory 存取数据。
4. GPU 将 GPU Memory 结果传回 Main Memory。

Grid、Block、Thread 概念：GPU 资源的抽象组织形式

但为了更好的利用 GPU 资源，提高并行度，CUDA 还要将这些 thread 加以优化组织，将能利用共有资源的线程组织到一个 thread block 中，同一 thread block 中的 thread 可以通过 share memory 共享数据，每个 thread block 最高可拥有 512 个线程。拥有同样维度同样 kernel 的 thread block 被组织成一个 Grid，而 CUDA 处理任务的最大单元便是 Grid 了。

如下图所示，利用 CUDA 进行编程时，一个 GRID 分为多个 Block，而一个 Block 可分为多个 Thread。

• 多个 Blocks 可以组成 1 维、2 维或者 3 维的 Grid。kernel 函数可以访问 Grid 内部标识 Block 的内置变量 BlockIDX，也可以访问表示 Block 维度的内置变量 BlockDim。
• 每个 Block 所能包含的 Thread 数量是有限制的，因为目前每个 Block 内的所有 Threads 都是在一个物理的处理器核中，并且共享了这个核有限的内存资源。当前的 GPU 中，每个 Block 最多能执行 1024 个 Thread。

Warp 概念：并发任务的调度单位

Warp 是 GPU 执行程序时的调度单位，目前 CUDA 的 Warp 大小为 32，同在一个 Warp 的线程，以不同数据资源执行相同的指令。

• Warp 调度单元与 Threads 的关系
  

上图反映了 Warp 作为调度单位的作用，每次 GPU 调度一个 Warp 里的 32 个线程执行同一条指令，其中各个线程对应的数据资源不同。

上图反映了 Warp 是如何排程的，即最终将 CUDA thread 映射到实际的物理硬件计算单元中执行（Thread => SP Block => SM Grid => GPU）。

一个 SM 只会执行一个 Block 里的 Warp，当该 Block 里 Warp 执行完后才会执行其他 Block 里的 Warp。进行划分时，最好保证每个 Block 里的 Warp 数量合理，那样一个 SM 可以交替执行里面的 Warp，从而提高效率。

此外，在分配 Block 时，要根据 GPU 的 SM 个数，分配出合理的 Block 数，让 GPU 的 SM 都利用起来，提利用率。分配时，也要考虑到同一个线程 Block 的资源问题，不要出现对应的资源不够。

CUDA threads 在执行时，可以访问多个 Memory Spaces，每个线程有自己的私有的 Local Memory。每个 Block 有一个 Shared Memory，Block 的所有线程都可以访问。

最后，所有线程都可以访问 Global Memory。不同的内存访问速度为：本地内存 > 共享内存 > 全局内存。

通过 cudaMalloc 函数分配的内存就是全局内存。核函数中用 __shared__ 修饰的变量就是共享内存。 核函数定义的变量使用的就是本地内存。

2007~2016：走向 HPC 和 DL/AI 智能

GPGPU 就绪后，GPU 生态大厦中其实还缺一块拼图 —— 应用。当传统上用于图形处理的 GPU 具备了高性能浮点运算能力，且能像通用 CPU 那样可编程，那它可以用来做什么呢？对于这个问题，中国的计算机科学家给出了答案 —— 科学计算与超级计算机。

2007 年，国防科大杨学军院士带领团队在国际计算机体系结构旗舰会议 ISCA 上发表题为《A 64-bit stream processor architecture for scientific applications》的学术论文，揭示了流处理器架构可大幅加速科学计算。

这项工作不仅为具备大量流处理器的 GPU 开辟了新应用场景，也为超级计算机架构设计开辟一条新技术路径，即采用 “CPU+GPU” 异构结构提升性能。

此后，GPU 便开始出现在超级计算机排行榜（Top500）上，并逐步成为主流：

• 2009 年 11 月，采用 Intel CPU+ATI GPU 异构结构的 “天河一号” 名列 Top500 排行榜第五；
• 2010 年 6 月，采用 Intel CPU+Nvidia GPU 异构结构的 “曙光 6000” 位列 Top500 排行榜第二；
• 2010 年 11 月，采用 Intel CPU+Nvidia GPU+FT-1000 异构结构的 “天河一号大江 A” 荣登 Top500 榜首，打破了长期以来美日霸榜的格局。

虽然彼时的英伟达 GPU 与 CUDA 已被证明是构建超级计算机的神器，但却尚未与这一轮 AI 浪潮直接关联起来。

普林斯顿大学李飞飞教授于 2009 年发布 ImageNet，随后为了推广 ImageNet，他们决定启动一项基于 ImageNet 的物体识别竞赛。2010 年第一届竞赛冠军识别错误率为 28%，2011 年第二届竞赛冠军错误率降到约 25%。

2012 年的第三届竞赛正是转折点 —— 加拿大多伦多大学的 Jeffrey Hinton 教授与其两位学生 Ilya Sutskever 和 Alex Krizhevsky 带着使用英伟达 GPU+CUDA 训练出来的深度神经网络 AlexNet 参加竞争，将错误率大幅降低了近 11 个百分点，高出第二名 41%。

算法、数据、算力在 2012 年的 ImageNet 竞赛上汇聚了，形成一条势不可挡的 “大江”，在学术界掀起了惊涛骇浪。而 NVIDIA 的 GPU 芯片与 CUDA 软件生态也正式搭上了 AI 快车，不断发展壮大。

NVIDIA 的 GPU 芯片和 CUDA 软件生态已在当前 AI 浪潮中占据算力主导地位并形成高度垄断。打破这种垄断格局成为全球的共识，Google、Meta 等企业都自研各自的 AI 处理器芯片。

今天，人们也常说这一轮 AI 浪潮离不开三个要素的汇聚，即算法、数据与算力。

2017~Now：引爆 AI 大模型时代

参考资料

1. https://positiveczp.github.io/%E7%BB%86%E8%AF%B4%E5%9B%BE%E5%BD%A2%E5%AD%A6%E6%B8%B2%E6%9F%93%E7%AE%A1%E7%BA%BF.pdf
2. https://mp.weixin.qq.com/s?__biz=MzUxNzQ5MTExNw==&mid=2247487976&idx=1&sn=d258826f5829b2225b93183248f3f893&chksm=f996012acee1883cd8a41eb8d57d6b6fee57357d5ffd184883a89a5b4524c796aec27e7840dc&cur_album_id=2538479717163761664&scene=189#wechat_redirect