GPU 绘制原理与架构

计算机图形学：基于GPU的绘制和计算

GPU 架构初探

一、绘制光栅图像的基本原理

1.1 物体形状的表示与处理

核心观点： 计算机通过将三维物体拆解成简单的、可处理的基本图元（主要是三角形），并利用齐次坐标与矩阵变换来确定它们在屏幕上的最终位置。

• 图元 (Primitive): 计算机图形学中用于构建复杂形状的最基本几何元素。

  • 核心图元是三角形： GPU 选择三角形作为渲染的基本单位，因为三点确定一个平面。这保证了图元的平坦性（Planarity），极大地简化了后续的光栅化计算和硬件设计。任何复杂多边形都可以被拆解为三角形。

• 网格曲面 (Mesh Surface): 通过大量的三角形拼接，来逼近和表示复杂三维物体的表面。这些模型由一系列顶点构成，每个顶点包含位置、法线、颜色、纹理坐标等属性。

• 坐标变换 (Coordinate Transformations): 一个顶点在最终显示到屏幕上前，需要经历一系列坐标空间的变换。这个过程本质上是向量与矩阵的乘法。

  • 变换流水线： 局部坐标系 (Model Space) → 世界坐标系 (World Space) → 观察坐标系 (View Space) → 投影坐标系 (Projection Space) → 屏幕坐标系 (Screen Space)。

• 齐次坐标 (Homogeneous Coordinates):

  • 核心思想： 使用 $N+1$ 维向量来表示 $N$ 维空间中的点或向量。例如，用四维向量 $(x,y,z,w)$ 表示三维空间。
  • 关键特性：
    • 当 $w=1$ 时，表示一个点 $(x,y,z)$。
    • 当 $w=0$ 时，表示一个方向向量 $(x,y,z)$。
  • 核心优势： 统一了平移、旋转、缩放等所有几何变换，使其都可以通过单一的矩阵乘法来完成，简化了数学表达和硬件实现。

• 裁剪 (Clipping) 与 剔除 (Culling):

  • 裁剪 (Clipping): 在投影变换后，由硬件自动执行，用于切掉部分超出视锥体的图元。
  • 剔除 (Culling): 通常指在流水线早期完整移除某些图元，例如背面剔除 (Back-face Culling) 会移除那些背对摄像机的三角形，以提升效率。

1.2 光栅化与像素生成

核心观点： 光栅化是将矢量化的几何图元（三角形）转换为像素集合（片元）的过程，随后经过一系列测试和计算，最终确定每个像素的颜色。

• 光栅图像 (Raster Image): 屏幕是由一个二维的像素矩阵构成的，光栅图像（或称位图）就是为这个矩阵中的每一个像素精确定义颜色值。

• 光栅化 (Rasterization):

  • 核心任务： 判断一个三角形覆盖了屏幕上的哪些像素中心点。
  • 生成片元 (Fragment): 对于每个被三角形覆盖的像素位置，生成一个片元。片元可以理解为一个“潜在的像素”，它携带了从三角形顶点插值而来的各种属性（如颜色、深度值、纹理坐标等）。

• 片元处理流水线 (Fragment Processing): 生成的片元需要经过一系列操作才能最终写入屏幕。

  • 着色 (Shading): 计算片元的最终颜色。这是通过片元着色器完成的，它会利用片元的属性（如法线、纹理坐标）和光照模型（如 Phong 模型）来进行计算。
  • 深度测试 (Depth Test): 使用 Z-Buffer 算法，比较当前片元的深度值与 framebuffer 中已存储的深度值。如果当前片元更远（被遮挡），则被丢弃。这是实现三维场景遮挡关系的关键。
  • 模板测试 (Stencil Test): 类似深度测试，但依据模板缓冲区的值来决定是否丢弃片元，常用于实现复杂的遮罩效果（如镜面、轮廓光）。
  • 颜色混合 (Color Blending): 将通过测试的片元颜色与 framebuffer 中已有的颜色按照特定规则进行混合，用于实现半透明效果。

二、初探 GPU 架构

2.1 绘制流水线与流式编程模型

核心观点： GPU 专为高吞吐量的并行计算而设计，其架构完美契合流式编程模型，这使其在图形渲染等领域远超 CPU。

• GPU vs. CPU 设计哲学:

  • CPU: 为低延迟 (Low Latency) 设计，擅长复杂的逻辑控制和串行任务。
  • GPU: 为高吞吐量 (High Throughput) 设计，拥有大量计算核心，擅长处理大规模、独立的并行数据。

• 流式编程模型 (Stream Programming Model): GPU 计算的核心思想。

  • 流 (Stream): 一系列类型相同的数据元素的有序集合（如顶点流、像素流）。
  • 核 (Kernel): 一个函数，它会被独立地应用到流中的每一个元素上。
  • 核心特性： 流中各元素的计算相互独立。这使得计算可以被大规模并行化，并且硬件可以预先优化数据访存，从而有效隐藏内存延迟。

• 多层次并行机制 (Multi-level Parallelism):

  • 任务级并行: 流水线的不同阶段（如顶点处理和片元处理）可以同时处理不同的数据。
  • 数据级并行: 同一阶段（一个 Kernel）可以同时处理流中的多个元素。
  • 指令级并行: 处理单个元素时，没有数据依赖的指令可以并行执行。

2.2 从固定功能到可编程流水线

核心观点： GPU 架构从早期功能写死的固定功能流水线，演变成了高度灵活的可编程流水线，将渲染过程中的关键步骤开放给开发者自定义，催生了现代计算机图形学的巨大飞跃。

• 固定功能流水线 (Fixed-Function Pipeline):

  • 特点： 早期的 GPU 将渲染流水线的每个阶段（如坐标变换、光照计算）都用专用硬件实现，功能固定，开发者只能通过 图形 API (Graphics API) 调整预设参数。
  • 代表 API 版本： 早期的 OpenGL 和 Direct3D (如 Direct3D 7 及以前)。

• 可编程流水线的诞生:

  • 里程碑： 2001 年，NVIDIA GeForce 3 显卡发布，首次支持可编程着色器。Direct3D 8 和后续的 OpenGL 2.0 正式将此概念引入 API。
  • 核心变革： 允许开发者编写自定义程序（Shader），来替代流水线中的某些固定功能模块。

• 核心可编程阶段:

  • 顶点着色器 (Vertex Shader): 针对每个顶点执行的程序。开发者可以完全控制顶点的位置、颜色等属性，从而实现复杂的模型动画、变形、粒子效果等。
  • 片段着色器 (Fragment Shader / Pixel Shader): 针对每个片元执行的程序。开发者可以实现任意复杂的光照模型、材质效果和后处理特效，精细控制最终像素的颜色。

• 最终演进: 随着技术成熟，固定功能流水线被完全移除（以 Direct3D 10 和 OpenGL 3.1 为标志），可编程流水线成为现代 GPU 的标准架构。这一变革不仅极大地丰富了图形表现力，也为 GPGPU（通用 GPU 计算） 的发展奠定了基础。

顶点/像素着色器、统一架构与 SIMT 模型

一、着色器编程入门 (Introduction to Shader Programming)

核心观点： 开发者通过 C++ 等高级语言调用图形 API (Graphics API)，将顶点数据、纹理等资源和编译好的着色器程序 (Shader) 提交给 GPU。着色器是运行在 GPU 上的小型程序，用于控制渲染流水线的可编程阶段。

1.1 高级着色器语言：GLSL vs. HLSL

为了取代早期类似汇编的复杂指令，现代图形开发普遍使用高级着色语言。其中最主流的是 GLSL 和 HLSL。

特性对比	GLSL (OpenGL Shading Language)	HLSL (High-Level Shading Language)
主要用途	OpenGL, Vulkan	Direct3D (Windows, Xbox)
版本控制	文件内声明 (#version 330 core)	编译时通过参数指定 (Shader Model)
数据流定义	使用 in, out, uniform 等存储限定符	使用语义 (Semantics) 附加在变量后 (如 : SV_POSITION)
入口点函数	统一命名为 void main()	可自定义多个不同名称的入口点函数
资源绑定	使用 layout(binding = ...) 进行显式布局	使用 register(b0, t0, ...) 显式绑定到虚拟寄存器
Vulkan 特性	需预编译成二进制中间格式 SPIR-V	-

1.2 核心概念：资源绑定与描述符

核心观点： 为了让着色器能访问纹理、缓冲区等外部资源，现代图形 API 引入了描述符 (Descriptor) 机制，它是一种标准化的资源“指针”，解耦了着色器代码与具体的资源内存。

• 描述符 (Descriptor): 一块包含访问资源所需全部信息的数据结构。例如，一个纹理的描述符会包含其内存地址、尺寸、格式等信息。
• API 的解决方案：
  • Vulkan: 描述符集 (Descriptor Set): 将资源描述符分组管理。开发者通常根据资源的更新频率（例如，每帧更新的、每个物体更新的）来组织描述符集，便于驱动进行优化。
  • Direct3D 12: 描述符堆 (Descriptor Heap): 一块连续的内存，用于存放大量的描述符。开发者拥有更高的控制权，可以直接管理这块内存。

二、GPU 架构的演进

2.1 从分离到统一：统一着色器架构

核心观点： 现代 GPU 采用统一着色器架构，将计算资源整合为统一的处理器池，可以根据实时需求灵活处理顶点、像素或其他任何着色任务，极大地提高了硬件利用率。

• 旧的“分离式架构”:

  • 设计: 拥有两种不同的、专门化的处理单元：一组用于顶点着色，另一组用于像素着色。
  • 缺陷: 资源无法动态调配。当场景顶点负载高而像素负载低时，像素处理单元闲置；反之亦然。这导致了严重的硬件资源浪费。

• 现代的“统一着色器架构 (Unified Shader Architecture)”:

  • 设计: 将所有处理单元设计成同一种通用的流式处理器 (Streaming Processor, SP)。
  • 优势: 任何一个 SP 都可以执行任何类型的着色器代码（顶点、像素、几何等）。GPU 的调度器会根据当前任务负载，动态地将任务分配给空闲的 SP，实现了负载均衡和效率最大化。
  • 推动力: Direct3D 10 API 的标准要求是推动这一架构变革的重要因素。

2.2 并行计算模型：VLIW, SIMD, 与 SIMT

核心观点： GPU 的并行计算模型从依赖编译器静态调度的 VLIW 演进到了由硬件动态调度的 SIMT 模型，从而更高效地处理复杂的、带有分支的现代着色器程序。

• SIMD (Single Instruction, Multiple Data): 单指令，多数据。这是所有并行处理器的基础。一条指令可以同时在多个数据上执行。例如，a = b + c 这条指令可以同时对 16 组不同的 b 和 c 进行求和。

• VLIW + SIMD (早期 AMD TeraScale 架构):

  • VLIW (超长指令字): 编译器负责将多条没有相互依赖的简单指令静态打包成一条“超长指令”，让硬件在一个周期内同时执行。
  • 缺点: 严重依赖编译器的优化能力。当着色器逻辑复杂、分支多时，编译器很难找到足够多的独立指令来填满一个 VLIW 包，导致大量硬件计算单元被闲置，执行效率低下。

• SIMT + SIMD (现代 NVIDIA 和 AMD GCN 之后架构):

  • SIMT (单指令，多线程): 硬件负责调度。它将大量线程（例如 32 或 64 个）打包成一个执行单元（Warp 或 Wavefront）。硬件每次只发射一条指令，这个指令会被包内的所有活动线程同时执行。
  • 优势: 灵活性和效率极高。硬件调度器可以在运行时动态处理分支（通过禁用不满足分支条件的线程）和隐藏内存延迟（当一个 Warp 等待内存时，调度器可以切换到另一个就绪的 Warp 继续执行）。

2.3 现代 GPU 并行编程抽象

核心观点： 尽管底层硬件以 Warp/Wavefront 为单位进行调度，但上层编程模型（如 CUDA）提供了更易于管理的抽象层次，让开发者能够组织和控制海量的并行任务。

以下是 NVIDIA (CUDA) 和 AMD 描述相似概念时使用的不同术语：

抽象层级	NVIDIA / CUDA 术语	AMD 术语	说 明
整个计算任务	网格 (Grid)	网格 (Grid)	对应一个完整的 Kernel 调用，例如处理一张图像的所有像素。
线程分组	线程块 (Thread Block)	工作组 (Work Group)	一组被调度到同一个计算单元（SM/CU）上执行的线程，可以共享该单元的局部内存。
硬件调度单元	Warp	Wavefront	硬件并行执行的基本单位，通常为 32 或 64 个线程，它们严格按照相同顺序执行相同指令。
最小工作单位	CUDA 线程 (Thread)	工作项 (Work Item)	执行计算的最小实体，通常对应一个数据元素（如一个顶点或一个像素）。

描述符绑定硬件实现差异

Descriptors are hard

核心问题

Vulkan 的一大复杂性来源是它需要支持种类极其繁多的硬件，而不同厂商在描述符绑定 (Descriptor Binding) 这一核心功能上的硬件实现方式千差万别。为了理解 Vulkan API 设计背后的考量，我们需要先了解底层硬件的几种主流实现思路。

四种主流硬件绑定方法

硬件上的描述符绑定机制大致可以分为以下四大类，每种都有其独特的优缺点。

1. 直接访问 (Direct Access - D)

• 工作方式: 着色器 (Shader) 直接将完整的描述符信息（可以理解为一个“原始指针”）传递给访存指令。
• 优点: 极其灵活。因为描述符可以存放在内存的任何位置。
• 缺点: 效率较低。每次访问资源都需要在着色器中传递整个描述符，数据传输量大。

2. 描述符缓冲区 (Descriptor Buffers - B)

• 工作方式: 描述符被统一存储在一个或多个缓冲区 (Buffer) 中。着色器只需要知道这个缓冲区的基地址和一个指向目标描述符的偏移量 (Offset) 即可。
• 优点: 比“直接访问”更高效，因为着色器只需传递一个较小的偏移量。更换绑定的缓冲区代价也相对较低。
• 缺点: 硬件需要先从描述符缓冲区中读取描述符，然后再进行真正的资源访问，多了一步间接操作。

3. 描述符堆 (Descriptor Heaps - H)

• 工作方式: 所有同类型的描述符都存放在一个全局的表或堆 (Global Table/Heap) 中。着色器通过一个索引 (Index) 来访问。
• 优点: 着色器传输数据量最小，效率最高，因为只需要一个简单的索引。
• 缺点: 非常不灵活，且更换代价极高。更改全局描述符堆通常会导致整个 GPU 停顿 (GPU stall) 并清空缓存。

4. 固定硬件绑定 (Fixed HW Bindings - F)

• 工作方式: 这是最传统的方式。资源被绑定到数量有限的、固定的硬件槽位 (Slot) 上，通常通过命令流直接设置寄存器来完成。
• 优点: 硬件实现简单。
• 缺点: 灵活性最差，不符合现代“Bindless”的设计趋势。在现代硬件上，这种方式通常只用于渲染目标 (Render Target)、顶点/索引缓冲区等固定功能的部件。

---

主流硬件厂商方案对比

下表展示了当今主流 GPU 厂商针对不同资源类型所采用的绑定方法：

硬件厂商	纹理 (Textures)	图像 (Images)	采样器 (Samplers)	边界颜色	类型化缓冲区	UBOs	SSBOs
NVIDIA (Kepler+)	H	H	H		H	D/F	D
AMD	D	D	D	H	D	D	D
Intel (Skylake+)	H	H	H		H	H/D/F	H/D
Intel (pre-Skylake)	F	F	F		F	D/F	F
Arm (Valhall+)	B	B	B		B	B/D/F	B/D
Arm (pre-Valhall)	F	F	F		F	D/F	D
Qualcomm (a5xx+)	B	B	B		B	B	B
Broadcom (vc5)	D	D	D		D	D	D

注意: Intel pre-Skylake 架构虽然被标记为固定绑定 (F)，但它实际上是一种更灵活的、带有间接层的类堆模型（Binding Table），为后续架构的演进提供了基础。

几何着色器、曲面细分与计算着色器

一、几何着色器 (Geometry Shader, GS)

核心观点： 几何着色器是渲染管线中的一个可选阶段，它弥补了顶点着色器无法处理完整图元的缺陷，允许在 GPU 上动态地创建、销毁或修改几何体。

1.1 核心功能与应用场景

• 解决的问题：

  • 顶点着色器 (VS) 的局限： VS 每次只能处理单个顶点，无法访问同一图元中的其它顶点信息，也无法增删顶点。
  • GS 的突破： GS 的输入是完整的图元（如一个点、一条线或一个三角形），使其能够进行更复杂的几何操作。

• 核心能力：

  1. 访问图元所有顶点： 可以基于整个图元的形状进行计算（如计算法线、判断朝向）。
  2. 增删图元： 可以不输出任何图元（实现剔除），也可以输出一个或多个新图元。
  3. 改变图元类型： 可以输入一个点，输出一个四边形（由两个三角形 strip 构成），常用于公告板 (billboard) 或粒子效果。

• 典型应用：动态曲线细分

  • 为了绘制平滑的贝塞尔曲线，可以只向 GPU 传入 4 个控制点。
  • GS 接收这 4 个点，根据其到摄像机的距离动态计算一个合适的细分等级。
  • 在着色器内部循环，利用贝塞尔参数方程生成一系列细分的顶点，最后输出一条由多个短线段构成的 line_strip。
  • 三次贝塞尔曲线参数方程: $B(t)=(1-t{)}^3{P}_0+3(1-t{)}^{2}t{P}_1+3(1-t)t^2{P}_2+t^3{P}_3,t\in [0,1]$

1.2 严重的设计局限与性能问题

核心观点： 尽管功能强大，但几何着色器因其动态可变的输出数量，在现代 GPU 宽 SIMD 架构上存在严重的性能瓶颈，应谨慎使用或避免使用。

• 并行效率低下： GPU 以 Warp/Wavefront (一组 32 或 64 个线程) 为单位执行。如果组内一个 GS 线程输出 100 个顶点，而其它线程只输出 2 个，那么整个组必须等待最慢的那个线程执行完毕，导致大量计算单元闲置和“线程发散”(Divergence)。
• 内存与带宽压力： GPU 需要为 GS 的输出预留最坏情况下的内存空间，这可能导致缓冲区被放置在较慢的显存中，带来高延迟和带宽瓶頸。
• 破坏顶点缓存： GS 生成的每个顶点都被视为全新的，无法利用顶点后处理缓存 (Post Transform Cache) 来复用相同的顶点，增加了额外的计算负担。
• 适用建议： 仅用于处理整个图元的少量任务，如轮廓线检测，不应用于大规模的几何体生成或曲面细分。

二、曲面细分阶段 (Tessellation Stage)

核心观点： 曲面细分是 Direct3D 11 / OpenGL 4.0 引入的、专用于高效生成平滑表面细节的硬件加速阶段。它通过将低多边形的控制面片 (Patch) 细分为大量微小三角形，实现了高质量的动态细节层次 (LOD)。

2.1 为何需要曲面细分？

几何着色器不适合做大规模、规则的表面细分。曲面细分阶段通过更结构化、更适合并行处理的方式解决了这一问题，避免了 GS 的性能陷阱。

2.2 曲面细分的三个核心阶段

曲面细分由三个协同工作的单元组成：

1. 外壳/控制着色器 (Hull/Tessellation Control Shader - HS/TCS):
   • 类型： 可编程。
   • 任务： 接收输入的控制点组成的面片 (Patch)，并决定“如何细分”。它计算并输出细分因子 (Tessellation Factors)，告诉下一阶段要生成多少新的顶点。
2. 镶嵌器 (Tessellator):
   • 类型： 固定功能硬件单元。
   • 任务： 接收 HS/TCS 输出的细分因子，并据此生成新顶点的参数化坐标 (如 UV 坐标)。它只负责生成细分的 “模式”或“拓扑结构”，不计算实际的顶点位置。
3. 域/评估着色器 (Domain/Tessellation Evaluation Shader - DS/TES):
   • 类型： 可编程。
   • 任务： 接收来自镶嵌器的一个参数化坐标，并访问整个面片的所有原始控制点。它的核心工作是根据参数方程（如贝塞尔、NURBS）计算出新顶点的最终位置和其它属性。

2.3 相对优势与局限

• 优势：
  • 高度并行友好： 其工作模式（特别是 DS/TES）与顶点着色器类似，非常适合 GPU 的并行架构。
  • 硬件加速： 镶嵌器是专用的高速硬件，效率极高。
  • 带宽高效： 只需要传输低分辨率的控制点，大大节省了内存带宽。
• 局限：
  • 灵活性较低： 仅适用于对表面进行规则化的细分，无法像 GS 那样进行任意的几何创造。

三、通用计算的演进：从流输出到计算着色器

3.1 流输出 / 变换反馈 (Stream Output / Transform Feedback)

核心观点： 该机制允许将几何处理阶段 (VS 或 GS) 的输出结果直接捕获到 GPU 缓冲区中，而无需经过光栅化，是 GPGPU 的早期雏形。

• 解决的问题： 在此之前，如果想在 GPU 上复用顶点变换后的数据（例如用于物理模拟或多趟渲染），唯一的办法是将其“绘制”到纹理中，流程繁琐且低效。
• 核心功能： 在几何处理和光栅化之间提供一个“旁路”，将顶点数据流直接写入 Buffer。
• 局限性： 仍然受限于图形渲染管线的框架和严格的输出顺序要求，在某些架构上性能开销较大。Vulkan 最初甚至没有包含此功能。

3.2 计算着色器 (Compute Shader, CS)

核心观点： 计算着色器是一个独立于传统图形管线的、为通用并行计算 (GPGPU) 而设计的强大工具。

• 核心特性：

  • 独立管线： 它不属于 "VS → GS → PS" 的流程，通过 Dispatch 命令而非 Draw 命令启动。
  • 通用内存访问： 可以自由地对 GPU 内存（缓冲区、纹理）进行读写，不受图形管线输入/输出的限制。
  • 灵活的线程组织： 开发者可以定义一个三维的计算网格，将任务划分为线程组 (Thread Group) / 工作组 (Work Group)，并进一步在组内进行同步和数据共享。

• 并行编程抽象（NVIDIA vs. AMD）:

抽象层级	NVIDIA / CUDA 术语	AMD 术语	说 明
整个计算任务	网格 (Grid)	网格 (Grid)	对应一次 Dispatch 调用的全部工作。
线程分组	线程块 (Thread Block)	工作组 (Work Group)	一组被调度到同一个计算单元 (SM/CU) 上的线程。
硬件调度单元	Warp	Wavefront	硬件并行执行的基本单位，通常为 32 或 64 个线程。
最小工作单位	CUDA 线程 (Thread)	工作项 (Work Item)	执行计算的最小逻辑实体。

---

GPU 微架构：硬件执行模型，调度与依赖管理

1. 从软件抽象到硬件执行

GPU 采用 单程序多数据（SPMD） 的形态，即一段相同的程序（Shader/Kernel）被应用于海量独立的数据元素（如顶点、像素）。为了高效管理，GPU 在软件抽象和硬件执行上采用了分层结构。

1.1 层次化结构映射

• 逻辑抽象（软件层）：
  • 网格（Grid）： 对应一次 Kernel 启动或 Draw Call，包含所有并行任务。
  • 线程块（Thread Block）： Grid 被分解为多个 Block，Block 是逻辑上的基本组织单位。
  • 线程（Thread）： 处理单个元素的最小单位。
• 物理执行（硬件层）：
  • GPC & SM： GPU 由多个 GPC（图形处理集群）组成，GPC 包含多个 SM（流式多处理器） 。
  • 驻留关系： 一个 Thread Block 驻留在单个 SM 上，共享该 SM 的资源（如 Shared Memory）。
  • 线程束（Warp）： 硬件调度的 最小并行单位 。SM 将 Block 内的线程进一步组织为 Warp（通常 32 个线程）。

1.2 调度机制

• Warp 调度： SM 中的 Warp 调度器在每个时钟周期选择 Block 内的某个 Warp 执行。
• SIMT 执行： Warp 内的线程以锁步（Lock-step）方式在 SIMD 单元上执行相同指令。
• 抽象与实现的差异： 开发者编写的是针对单个 Thread 的逻辑，而编译器和硬件负责将其聚合为 Warp 级别的指令流。

1.3 术语对照表（NVIDIA vs AMD）

概念	NVIDIA 术语	AMD 术语	定义
逻辑整体	Grid	Grid	并行任务全集
逻辑分组	Thread Block	Work Group	逻辑上的协作组，驻留同一计算单元
硬件调度单位	Warp	Wave / Wavefront	硬件执行的最小 SIMD 集合
最小处理单元	CUDA Thread	Work Item	单个元素的处理逻辑

---

2. SIMD 硬件与 SIMT 抽象

现代 GPU 普遍基于 SIMD（单指令多数据） 硬件构建，但在编程模型上呈现为 SIMT（单指令多线程） 。

2.1 SIMT 核心机制

• 编程体验： 程序员只需关注单个数据元素的标量逻辑（Scalar），无需像 CPU SSE/AVX 那样编写显式的向量代码。
• 硬件实现： 硬件将多个线程（如 32 个）打包成 Warp，共用一个取指单元。
• 分支分歧（Branch Divergence）：
  • 当遇到 if-else 时，硬件通过 执行掩码（Masking） 处理。
  • 串行化执行： 先执行 if 分支（屏蔽 else 线程），再执行 else 分支（屏蔽 if 线程）。这会导致 SIMD 效率下降。

2.2 架构演进案例：AMD 从 VLIW 到 SIMT

AMD 的架构演变清晰地展示了从“编译器静态调度”向“硬件动态调度”的转型。

A. TeraScale 架构（VLIW 模型）

• 核心设计： 结合 SIMD 与 VLIW（超长指令字） 。每个 SPU 包含 5 个 SP（VLIW5），通过编译器将 5 条独立指令打包成一条超长指令并行执行。
• 工作流：
  • 输入：$v_x^{\prime }=M_{11}{v}_x+M_{12}{v}_y+M_{13}{v}_z+M_{14}{v}_w$
  • 编译器打包： 必须静态分析指令依赖，填满 VLIW 的槽位（Slot）。
  • 缺陷： 对于复杂逻辑，编译器难以找到足够的独立指令填满槽位，导致大量 空操作（NOP） ，硬件利用率低。

B. GCN 架构（纯 SIMT 模型）

• 核心设计： 放弃 VLIW，转向纯 SIMT。基本单元 CU 包含向量 ALU（Vector ALU）和 标量 ALU（Scalar ALU） 。
• 动态调度： 硬件在运行时维护指令 buffer，动态发射指令，不再依赖编译器复杂的静态打包。
• 标量/向量分离：
  • Vector ALU： 处理每个线程不同的数据（如 $M_{11}{v}_x$）。
  • Scalar ALU： 处理 Wavefront 中所有线程相同的指令（如循环计数器），结果广播给全组，极大提升效率。

---

3. 数据冒险与依赖管理

GPU 通过在同一核心上交替执行多个 Warp 来 隐藏延迟 。为了保证执行正确性，必须处理 数据冒险（Data Hazards） 。

3.1 数据冒险类型

假设指令序列为 $I1$（写 $R$） $\to$ $I2$（读/写 $R$）：

1. RAW（写后读）： $I2$ 需要 $I1$ 的结果。这是 真依赖（True Dependency） ，必须等待 $I1$ 完成。
2. WAR（读后写）： $I2$ 写入 $I1$ 的输入寄存器。需防止 $I2$ 过早覆盖旧值（反依赖）。
3. WAW（写后写）： $I2$ 覆盖 $I1$ 的输出。需保证写入顺序（输出依赖）。

3.2 记分牌机制（Scoreboard）

GPU 使用 记分牌 硬件动态追踪指令状态，决定 Warp 是否“逻辑就绪”。

• 逻辑就绪条件：
  1. 激活： Warp 未结束且未被挂起。
  2. 同步： 无未满足的 Barrier。
  3. 数据依赖： 记分牌 确认所有操作数已就绪（RAW 解决）。
• 双记分牌策略（示例）：
  • 记分牌 1（生产者）： 跟踪未完成的写入，阻止 RAW 和 WAW。
  • 记分牌 2（消费者）： 跟踪正在进行的读取，阻止 WAR。

3.3 演进：编译器辅助管理

随着并发量增加，纯硬件记分牌的面积与功耗开销过大。现代架构（如 NVIDIA）趋向于 软硬协同 ：

• 机制： 编译器在指令中编码依赖信息（如设置等待位、计数器）。
• 优势： 简化硬件电路，降低功耗，将部分依赖分析工作前置到编译期。

寄存器与存储层次结构

1. 寄存器文件：并行度的首要瓶颈

在 SIMT（单指令多线程）模型下，寄存器不仅是速度最快的存储层级，更是决定 GPU 并行度上限的 核心资源 。由于大量活跃线程（Warp/Wavefront）需要同时驻留，指令对寄存器数量和带宽的需求被成倍放大。

1.1 核心概念与约束

• 寄存器文件（Register File, RF） ：片上高速存储，用于保存线程的私有状态（中间值、结果）。RF 的容量直接限制了单个 SM/CU 上能同时运行的 活跃线程数 。
• 占用率（Occupancy） ：指 SM/CU 上实际驻留的 Warp 数量与最大可能数量的比率。
  • 高占用率的意义 ：利用多组 Warp 轮转执行，以掩盖高昂的内存访问延迟（Latency Hiding）。
  • 权衡 ：为了追求极致的单线程性能而分配过多寄存器，会导致驻留 Warp 减少，降低占用率，从而削弱隐藏延迟的能力。

1.2 寄存器分配与溢出

• 寄存器分配（Register Allocation） ：编译器将局部变量映射到硬件寄存器的过程（NP-难问题）。
  • 动态分配 ：如 AMD RDNA4 允许根据负载动态调整寄存器量，光追场景下甚至可在生命周期内调整。
• 寄存器溢出（Register Spilling） ：当寄存器需求（Register Pressure）超过硬件限制时，编译器将部分变量移至 本地内存（Local Memory） 。
  • 策略 ：有时为了维持较高的占用率，编译器会主动选择溢出部分变量。
  • 代价 ：由于 Local Memory 通常位于显存（VRAM）并经由缓存访问，溢出会带来巨大的延迟和带宽开销。

1.3 硬件设计挑战与优化

• 分体式结构（Banked Structure） ：为解决多端口读写的物理限制，将寄存器划分为多个 存储体（Bank） 。
  • 存储体冲突（Bank Conflict） ：当一条指令的多个操作数映射到同一 Bank 时，读写被迫串行化，导致延迟增加。
• 层次化设计（Hierarchical RF） ：为缓解主寄存器压力，部分架构（如研究中的模型）引入中间缓存层：
  • LRF (Last Result File) ：保存极短生命周期的临时值。
  • ORF (Operand Register File) ：由软件（编译器）管理的缓存，利用时序局部性减少对主寄存器（MRF）的访问。

2. 谓词与标量寄存器：控制流与数据的优化

2.1 谓词寄存器（PRF）：解决分支分歧

当 Warp 内线程因条件判断产生 分支分歧（Branch Divergence） 时，GPU 利用 谓词寄存器文件（Predicate Register File） 进行管理，而非传统的 CPU 跳转。

• 执行掩码（Execution Mask） ：PRF 本质是一个位掩码，每一位对应一个线程。
• 锁步执行 ：
  • 若谓词位为 true ：线程执行指令并写入结果。
  • 若谓词位为 false ：线程执行 NOP（空操作），不写入结果。
• 优势 ：编译器将控制流转换为带掩码的顺序指令流，避免了上下文切换开销。

2.2 标量寄存器（sGPR）：消除数据冗余

针对 Warp 内所有线程共享的 统一数据 （Uniform Data，如循环计数器、常量地址），AMD GCN/RDNA 等架构引入了标量寄存器。

• 机制 ：数据只存储一份，所有线程通过广播读取。
• 收益 ：
  1. 降低向量寄存器（vGPR）压力，提升占用率。
  2. 利用 标量 ALU（SALU） 处理地址生成等逻辑，可与向量 ALU（VALU）并行发射，提升指令级并行度（ILP）。

3. GPU 缓存体系：带宽过滤器

GPU 缓存设计的核心目标并非 CPU 式的低延迟与强一致性，而是 吞吐量（Throughput） 。它充当 “带宽过滤器” ，旨在利用局部性减少对高能耗显存（DRAM）的访问事务。

3.1 层次结构与特性

• L1/纹理/常量缓存 ：针对特定访问模式特化。
  • 纹理缓存 ：优化 2D 空间局部性。
  • 常量缓存 ：优化广播与只读共享。
• 片上共享内存（Shared Memory / LDS） ：
  • 显式管理 ：由开发者编程控制（如 CUDA __shared__）。
  • 用途 ：Warp/Block 内的高效数据交换与复用。
  • 风险 ：需处理 Bank Conflict 和同步（Barrier）。
• L2 缓存 ：跨 SM/CU 的数据枢纽，主要负责过滤片外流量。

3.2 可见性与一致性模型

GPU 采用 弱一致性 模型。

• Warp 内 ：天然一致。
• 跨线程/跨 Block ：缓存命中不保证其他核心可见。必须显式使用 同步原语（Barrier/Atomic） 来保证内存顺序和可见性。

3.3 显存（VRAM）访问优化

即使拥有高带宽，实际性能仍受限于访问模式：

• 合并访问（Memory Coalescing） ：SIMT 单元尝试将同一 Warp 内相邻线程的访存请求合并为极少量的内存事务（Transaction）。
  • 未合并（Uncoalesced） ：稀疏、随机或未对齐的访问会导致事务数量激增，带宽利用率大幅下降。

4. 渲染架构：TBDR 与 Tiling

为了应对带宽瓶颈，图形管线在几何与光栅化阶段引入了分块处理思想。

4.1 移动端：TBDR (Tile-Based Deferred Rendering)

移动 GPU（如 Adreno, Mali, Apple）为降低功耗，广泛采用 TBDR。

• 核心流程 ：
  1. Binning ：几何阶段结束后，不立即光栅化，而是将图元分配到屏幕对应的图块（Tile）列表。
  2. On-Chip Rendering ：逐 Tile 将数据读入片上高速缓存，完成所有像素处理（光栅化、着色、混合）。
  3. 延迟写回 ：仅将最终颜色结果写回显存。
• 优势 ：将高频的深度/颜色读写（Depth/Color R/W）限制在片上，大幅减少 DRAM 带宽。
• 代价 ：增加了几何阶段的延迟（需等待全帧几何信息）和 Parameter Buffer 的开销。

4.2 桌面端：IMR 与 Tile Caching

桌面 GPU（NVIDIA, AMD）传统上使用 立即渲染模式（IMR） ，但也融合了分块思想。

• Tile Caching ：如 NVIDIA Maxwell 和 AMD Vega，在光栅化阶段利用片上 Bin Cache 提前剔除不可见像素，并优化 L2 缓存命中率。
• 区别 ：桌面端通常不进行全帧几何的延迟与分桶，而是侧重于利用分块来提升光栅化期间的缓存局部性。

总结：从管线到计算的演进

• 资源妥协 ：寄存器大小、带宽和延迟之间的权衡贯穿 GPU 架构设计。
• 局部性优化 ：从寄存器层级、缓存特化到 TBDR 架构，核心逻辑都是通过挖掘数据的时空局部性，将高频数据交互限制在离 ALU 最近、能耗最低的存储层级上。
• 趋势 ：由于传统几何管线的带宽效率瓶颈，业界正向 网格着色器（Mesh Shader） 演进，通过将处理粒度从“图元”提升到“微网格（Meshlet）”，以更可控的方式管理片上数据复用。

计算着色器（Compute Shader）与 GPU 并行编程模型

一、计算着色器的诞生与定位

1.1 从 GPGPU 到 Compute Shader 的演进

• 早期 GPGPU 开发者被迫将计算问题 伪装成图形任务 ——通过"画像素"来算数值，数据表示和内存访问都严重受限
• Direct3D 10 / OpenGL 3.0 引入的 流输出（Stream Output） 缓解了 CPU-GPU 数据搬运问题，但仍绑定在几何阶段的编程模型上
• 计算着色器 在 Direct3D 11（2009）和 OpenGL 4.3（2012）中正式引入，彻底脱离图形管线，具备三个关键能力：
  • 显式访问高速共享内存（LDS/groupshared） ：线程组内通信
  • 非连续写入与任意索引访问 ：不再受限于像素/顶点的规则输出
  • 原子操作（Atomic Operation） ：安全处理竞争写入

1.2 异步计算（Async Compute）

• AMD RDNA 架构中， 图形命令处理器（Graphics CP） 管理图形管线任务， 异步计算引擎（ACE） 管理计算着色器任务，两者 共享计算单元（CU）、L2 缓存、显存带宽
• 核心动机 ：图形管线天然存在 "资源气泡" ——某些阶段某类硬件空闲，异步计算可以"见缝插针"填满利用率
• 注意事项 ：
  • 计算任务可能 逐出图形任务的 L2 缓存数据
  • 长时间运行的计算任务会引发 上下文切换开销
  • 最适合 小粒度、低延迟 的任务（如剔除、物理模拟），而非大规模重度计算

1.3 编程模型：Dispatch 与线程网格

• 通过 Dispatch(GroupCountX, GroupCountY, GroupCountZ) 启动线程组网格
• 着色器用 [numthreads(X, Y, Z)] 声明每个线程组的维度
• 三个关键语义：
  • SV_GroupID ：线程组在网格中的 ID
  • SV_GroupThreadID ：线程在组内的 ID
  • SV_DispatchThreadID ：全局唯一线程 ID
• 核心思想 ：只要能把任务表达成 "可并行元素集合" ，就能用统一的线程网格承载

---

二、重构算法以适配 GPU：并行基数排序实例

2.1 两个核心并行原语

前缀和（Prefix Sum / Scan）

• 给定 $[a_0,a_1,a_2,\cdots ]$，输出 $[0,a_0,a_0+a_1,a_0+a_1+a_2,\cdots ]$（exclusive scan）
• 解决"并行输出位置"问题 ：每个线程要输出不同数量的数据时，对输出数量做一次前缀和，即得每个线程的起始写入偏移
• GPU 缺乏全局顺序执行模型，前缀和几乎是 唯一可行的并行位置确定方案

Scatter 与 Gather

• Gather（聚读） ：Output[i] = Input[Index[i]]——间接读取，GPU 处理较好（纹理采样即此模式）
• Scatter（散写） ：Output[Index[i]] = Value[i]——间接写入，是并行编程的 核心难点 ：
  • 多线程写同一位置 → 竞争
  • 写入地址分散 → 带宽利用率极低

2.2 为什么选择基数排序（Radix Sort）

• CPU 常用的快速排序依赖 递归 + 分支预测 ，在 GPU 的 SIMT 模型下效率极差
• 基数排序 不比较大小 ，而是 按位分桶 ：从低位到高位，每轮处理若干 bit（如 4 bit → 16 个桶），多轮"分桶 + 重排"后自然有序
• GPU 实现分 三个阶段 ：

阶段	名称	功能
1	Histogram（直方图统计）	统计每个桶有多少元素
2	Scan（前缀扫描）	计算每个桶的起始偏移量
3	Scatter（重排）	根据偏移量将元素移到新位置

2.3 优化策略一：合并内存访问（Coalesced Access）

• GPU 全局内存读写以 事务（Transaction） 为最小单位（如 32 字节一个 sector）
• 一个 32-wide wave 需要加载 32 个 float（128 字节）：
  • 地址连续 → 4 次 32B 事务，有效利用率 100%
  • 地址随机 → 最坏 32 次独立事务（传输 1024B，有效仅 128B），利用率 12.5%
• 核心策略 ：让 相邻线程读取相邻地址 ，哪怕多做 ALU 计算来组织地址也值得

实践示例 ：RadeonRays 的 BlockScan 函数让每个线程组的 256 个线程分 4 次迭代协同读取连续地址：

thread 0 → keys[0],   keys[256], keys[512], keys[768]
thread 1 → keys[1],   keys[257], keys[513], keys[769]
...
thread 255 → keys[255], keys[511], keys[767], keys[1023]

每次迭代中相邻线程访问相邻地址 → 合并为少量事务。

Scatter 阶段的困境 ：写入地址由数据值决定，天然随机不连续。高性能实现会 先在 LDS 内局部重排，凑成连续块后再统一写出 ——用片上带宽换片外带宽。

2.4 优化策略二：避免 LDS 存储体冲突（Bank Conflict）

• LDS 通常由 32 个独立存储体（Bank） 构成，每个周期每个 Bank 只能服务一次访问
• 理想：wave 内 32 个线程访问 32 个不同 Bank → 零冲突并行
• 冲突：多个线程命中同一 Bank → 串行化，延迟翻倍

代码中的转置技巧 ：

groupshared int lds_loads[4][256];  // [KEYS_PER_THREAD][GROUP_SIZE]
// 步骤2读取时：lds_loads[i][lidx]
// thread 0 → bank (256*0+0)%32 = 0
// thread 1 → bank (256*0+1)%32 = 1
// ...
// thread 31 → bank (256*0+31)%32 = 31

由于 GROUP_SIZE=256 是 32 的整数倍，列索引 lidx 直接映射到不同 Bank → 零冲突 。

注 ：计算着色器线程组内的线程按索引 连续打包 到 wave 中（像素着色器因 2×2 quad 规则则不同）。

2.5 优化策略三：算力与带宽的权衡

核心原则：GPU 上带宽通常比算力更稀缺

• RadeonRays 的第三阶段（Scatter）重新计算 了第一阶段的直方图，而非从显存读取中间结果
• 原因：重复计算的 ALU 代价 < 额外读写 VRAM 的带宽代价

反面案例——过度优化的陷阱 ：

步骤 1 向 LDS 写入时存在 Bank Conflict，但分析后发现：

• 写入 LDS 时的冲突 被 VRAM 读取延迟掩盖 ，消除它不会加速
• 简单 线性填充（Padding） 对"窄矩阵 + 宽 wave"的组合无效——冲突从"垂直"变成"对角线"
• XOR Swizzling 可以消除冲突，但需要额外 ALU 指令计算地址

$\text{swizzled\_col}=\text{col}\oplus (\text{row}\ll 3)$

工程结论 ：除非性能分析工具（如 Nsight Compute）确认带宽是瓶颈，否则 不要为了带宽优化而盲目增加算力需求 。

2.6 优化策略四：Wave Intrinsics 绕过 LDS

• 现代 GPU 提供 Wave Intrinsics （如 WavePrefixSum），利用 SIMD 单元内部的 交叉开关网络（Crossbar Network） 直接读取同一 wave 内其他线程的寄存器
• 完全绕过 LDS → 无 Bank Conflict、无同步屏障、节约 LDS 空间
• 最佳实践 ：
  • Wave 内部 的规约/扫描 → 优先用 intrinsics（寄存器到寄存器）
  • 跨 Wave（Inter-wave） 的数据交换 → 才使用 LDS
  • 只有数据量超过寄存器容量、或需跨 wave 同步时，才回退到 LDS

两种 BlockScan 实现对比 ：

特征	经典 Blelloch Scan	Wave Intrinsics 版本
通信介质	全部通过 LDS	Wave 内走寄存器，仅跨 wave 用 LDS
同步屏障数	多次 GroupMemoryBarrierWithGroupSync	显著减少
Bank Conflict 风险	需要精心设计	Wave 内完全消除

---

三、GPU、NPU 与张量核心的架构折中

3.1 GPU 通用计算的局限

GPU 的"通用性"建立在 一系列强假设 之上：

• 计算负载可拆分为 大量独立、结构相似 的工作项
• 工作项能在 统一执行模型 下推进
• 优化目标是 吞吐 而非单次延迟

任何偏离这些假设的场景 （不规则控制流、强串行依赖、少量数据）都会导致效率急剧下降，且 不会"优雅退化" 。

3.2 NPU / 专用加速器的路线（以 AMD XDNA 为例）

核心设计思想 ：主动缩小问题空间，换取能效与确定性

特征	GPU	AMD XDNA NPU
调度方式	运行时动态调度	编译时静态规划
缓存管理	硬件自动	显式编程（DDR → L2 → L1，DMA 驱动）
物理架构	统一计算单元	Tile 阵列（计算 Tile + 存储 Tile + 接口 Tile）
多任务	上下文切换	物理分区（列级隔离，保证 QoS）

优势 ：极致能效比、确定性延迟 局限 ：算法偏离预设结构时灵活性急剧衰减；微小算法变更可能需要重构编译器甚至重新流片

3.3 张量核心（Tensor Core）——通用与专用之间的折中

定位 ：在通用 GPU 内嵌入 领域专用加速单元 ，匹配时走高效专用路径，不匹配时回退通用路径

代表性硬件 ：

• NVIDIA Turing：Tensor Core + RT Core
• AMD RDNA2：Ray Accelerator
• AMD RDNA3：AI Acceleration

3.4 NVIDIA Volta 张量核心的工作方式

突破 SIMT 模型 ：传统 SIMT 中每个线程拥有独立寄存器状态；张量核心要求 整个 warp 协同持有一个矩阵分片（Fragment） ，数据以 不透明的跨线程形式 分布在 warp 的寄存器堆中。

执行 $16\times 16$ 矩阵乘加 $C=A\times B+C_{in}$：

• 32-wide warp 分为 8 组 ，每组 4 个连续线程负责一个 $4\times 8$ 的结果子块
• 每个子块的点积 $A_{4\times 16}\cdot B_{16\times 8}$ 被拆成 4 步 ，每步处理 4 列/行
• 广播/复用机制 ：读一次 $A$ 的 $4\times 4$ 子块，同时乘以 $B$ 的两个并排 $4\times 4$ 子块 → 一次读取完成两组输出的计算 → 减少寄存器读取和带宽压力

编程方式 ：

// 不再用数组下标直接访问矩阵元素
// 而是通过专用 intrinsic 操作 fragment
wmma::load_matrix_sync(a_frag, A_ptr, lda);
wmma::load_matrix_sync(b_frag, B_ptr, ldb);
wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);
wmma::store_matrix_sync(C_ptr, c_frag, ldc, ...);

以 指令通用性 和 更严格的数据布局约束 为代价，换取 极致吞吐量 ——GPU 从"通用并行处理器"向 领域专用架构（DSA） 的部分回归。

---

四、核心要点总结

主题	关键认知
Compute Shader	脱离图形管线，把并行控制权交给程序员
Scan（前缀和）	GPU 并行编程中确定写入位置的基础原语
Coalesced Access	相邻线程读写相邻地址，最大化内存事务效率
Bank Conflict	LDS 分体结构下的访问冲突，需显式规避或用 Wave Intrinsics 绕过
算力 vs 带宽	GPU 上重新计算往往比多一次 VRAM 读写更划算
Wave Intrinsics	现代 GPU 的寄存器直接通信，优先于 LDS 用于 wave 内操作
GPU vs NPU	通用性（可编程、灵活）与专用性（能效、确定性）的持续权衡
Tensor Core	将 warp 从"独立线程集合"升级为"协同矩阵处理单元"，以约束换吞吐

现代几何管线：渲染管线几何阶段的计算化与网格着色器

---

一、渲染管线几何阶段的计算化

1.1 图元子组（Primitive Subgroup）与批量处理

• 传统管线通过 索引缓冲区（Index Buffer） + 后变换缓存（Post-Transform Cache） 复用已着色顶点，减少重复计算
• AMD RDNA 架构的 NGG（Next Generation Geometry） 技术将这一过程显式化，流程为：
  1. 从索引缓冲区取出一个图元子集
  2. 去重顶点索引，构建 图元子组 ——包含唯一顶点列表 + 重建的局部索引
  3. 以图元子组为粒度，批量着色并送入光栅化
  4. 循环处理下一批，直到所有图元渲染完毕
• 核心思想 ：把"逐顶点顺序处理"变成"按批次并行处理"，一个图元子组对应一个线程组在双计算单元（Dual CU）上执行

1.2 顶点重用的理想与现实

• 理想目标 ：每个顶点仅出现在一个图元子组中，只被着色一次
• 现实约束 ：
  • 图元子组大小有限，无法保证全局唯一
  • 三维重建等来源的网格，索引排列往往不优化
  • meshoptimizer 等工具可离线重排索引，将共享顶点的图元聚拢，提高缓存命中率

1.3 传统管线的瓶颈

瓶颈类型	具体表现
索引缓冲区瓶颈	庞大输入数据受限于固定格式（16/32-bit 索引），所有顶点属性共用一个索引值
剔除粒度不灵活	顶点着色器在剔除发生前就已执行，大量不可见顶点白白消耗带宽和算力
控制力不足	程序员无法灵活控制调用时机、数据格式、LOD 选择

1.4 计算化的两条路径

• 纯软件方案 ：用计算着色器完全替代几何阶段
  • 典型案例：虚幻 5 的 Nanite ——手动维护任务队列，自由控制 LOD 遍历、聚类剔除，甚至用软件光栅化替代硬件光栅化
• 硬件 + API 方案 ：网格着色器 + 放大着色器（任务着色器） ，在保留硬件光栅化的前提下，将几何阶段计算化
• 共同核心 ：用 "计算式线程组协作模型" 替换传统的 "逐顶点/逐图元顺序执行模型"

---

二、网格着色器（Mesh Shader）

2.1 Meshlet：几何数据的基本处理单元

• Meshlet（微网格） = 小型、独立的几何数据包，包含：
  • 一组 非重复顶点 及其属性
  • 一组 局部图元索引 （引用上述顶点）
  • 可选的包围盒 / 法线锥等元数据（用于早期剔除）
• 典型规模：64 顶点 / 128 图元 （与硬件执行单元匹配）

2.2 编程模型

• 调度方式类似计算着色器：

graphicsCommandList->DispatchMesh(groupCountX, groupCountY, groupCountZ);

• 着色器内部用 [NumThreads(N, 1, 1)] 声明线程组大小，用 [OutputTopology("triangle")] 声明输出拓扑
• 关键 API 调用 ：SetMeshOutputCounts(vertCount, primCount) ——必须在写入输出前至少调用一次，告诉硬件本组将输出多少顶点和图元
• 线程组协作流程：
  1. 从自定义格式的缓冲区中 主动加载 meshlet 数据（绕过输入装配器）
  2. 组内线程并行处理顶点变换、图元索引解包
  3. 通过 共享内存（LDS） 实现顶点重用和组内通信
  4. 输出顶点属性数组 + 图元索引数组，交给后续光栅化

对比维度	传统管线	网格着色器
处理单元	单个顶点 / 单个图元	整个 meshlet（线程组协作）
输入方式	输入装配器解析固定格式索引	程序员自定义格式，主动抓取
顶点重用	依赖硬件后变换缓存	通过共享内存显式复用
控制权	硬件驱动	程序员完全控制

2.3 对比几何着色器（Geometry Shader）

几何着色器看起来功能相似（能增删改图元），但两者的架构差异是根本性的：

维度	几何着色器	网格着色器
粒度	单线程处理单个图元	线程组协作处理整个 meshlet
线程间通信	互相孤立，无法共享	通过 LDS 共享数据，天然支持顶点重用
输出空间	每线程独立预留缓冲（最坏情况：$32\times 4\text{KB}=128\text{KB}$）	整个线程组共享一个输出缓冲
占用率影响	单个 warp 可能占满 CU 全部片上存储 → 占用率极低	资源需求小数量级 → 可并发更多 wave
剔除时机	只能在图元级别剔除，顶点着色已完成	可配合放大着色器在 meshlet 级别 提前剔除
本质定位	对传统管线的 修补/扩展	对几何阶段的 重构

2.4 放大着色器 / 任务着色器（Amplification / Task Shader）

• 位于网格着色器 之前 的可选阶段，功能定位：粗粒度早期剔除 + 动态工作量生成

• 典型工作流：

  1. 每个线程检查一个 meshlet 的可见性（视锥体剔除、背面锥剔除等）
  2. 用 Wave 内置函数 统计可见 meshlet 数量：
     • WavePrefixCountBits(visible) → 确定每个可见 meshlet 在输出中的紧凑索引
     • WaveActiveCountBits(visible) → 得到总可见数量
  3. 将可见 meshlet 索引写入 payload（载荷） ——组内共享内存中的数据结构
  4. 调用 DispatchMesh(visibleCount, 1, 1, payload) 动态发射网格着色器线程组

• 其他应用场景 ：

  • 动态 LOD ：根据摄像机距离选择不同精度的 meshlet
  • 程序式几何生成 ：类似曲面细分中外壳着色器控制细分级别的角色
  • 需要精心设计 payload 数据结构，在两个阶段间高效传递实例化参数

2.5 网格着色器的局限与工程权衡

① 数据传递与调度开销

• AS → MS 的 payload 传递经过片上共享内存，延迟远高于寄存器通信
• 若 meshlet 粒度过小（<64 顶点），额外的 payload 带宽和硬件调度开销 可能抵消剔除收益

② 存储资源占用限制并行度

• MS 线程组必须在片上显式分配整个 meshlet 的输出空间（所有顶点属性 + 图元索引）
• 声明过大的 max_vertices / max_primitives → 单组占用过多 LDS/寄存器 → 降低占用率（Occupancy） → 无法隐藏访存延迟
• 工程准则 ：meshlet 规模必须与硬件执行单元严格匹配（如 64 顶点 / 128 图元）

③ 特定场景下的劣势

• 简单位移映射 / 高倍率细分 ：GPU 内置的 曲面细分固定功能硬件 在能效和速度上优于通用计算逻辑
• 用网格着色器实现等效功能需手动管理动态拆分、顶点生成、索引重排 → 开发维护成本高

---

三、硬件实现：AMD NGG 的两级着色器映射

AMD RDNA 在硬件层面只用 两个着色阶段 实现所有软件几何着色器：

硬件阶段	名称	职责
表面着色器	Surface Shader（Pre-Tessellation）	概念上对应外壳着色器
图元着色器	Primitive Shader	处理图元子组，输出图元 + 顶点信息

软件着色器到硬件阶段的映射：

软件组合	表面着色器	图元着色器
VS	—	VS
VS + HS + DS	VS + HS	DS
VS + GS	—	VS + GS
VS + HS + DS + GS	VS + HS	DS + GS
MS	—	MS
AS + MS	AS	MS

• 传统 VS 路径：几何引擎从索引缓冲区加载索引 → 构建图元子组 → 检查顶点重用 → 启动图元着色器线程
• 网格着色器路径 ：几何引擎启用 快速启动模式 ，绕过 硬件固定功能的顶点重用检查和图元子组构建（这些由程序员在着色器内手动处理）

---

四、GPU 工作图（GPU Work Graphs）

4.1 解决的核心问题

即便有了 AS + MS，渲染管线仍受制于 CPU 发起 → GPU 执行 的线性交互模式：

现有机制	局限
传统 Draw/Dispatch	每次调度都需 CPU 介入
ExecuteIndirect	GPU 只能填充预分配命令的参数，无法改变命令拓扑——本质是 "预设填空题"

两难困境 ：

• 频繁 CPU-GPU 往返 → 同步延迟
• 在一次 dispatch 中处理所有情况 → 编写 "超级着色器（Uber Shader）" → 寄存器需求激增 → 占用率暴跌

4.2 工作图的核心模型

D3D12 工作图在 GPU 内部构建一个由节点（Node）组成的 有向图 ，实现细粒度的 "生产者-消费者" 模型：

• 自驱动调度 ：正在执行的着色器（生产者）计算出局部结果后，将数据打包为 记录（Record） ，直接发射给后续的不同着色器节点（消费者）
• 片上数据流 ：新任务无需回传显存等待 CPU，直接进入硬件任务队列。数据极有可能在 L2 甚至片上缓存 中被消费者直接复用
• 按需分配 ：节点若不产生有效数据，后续节点被 静默跳过 ，算力不被浪费

4.3 适用性权衡

• 适合 ：高度动态的异构任务（不同材质调用不同着色器、光追次生射线等）
• 不适合 ：结构规整、大批量的同质化任务——此时传统 ExecuteIndirect 或简单计算着色器更优
• 新的权衡维度 ： 调度灵活性 vs. 数据局部性 ——动态图调度可能打破原本规整的访存模式

---

五、关键要点总结

传统管线                          现代几何管线
Input Assembler              Amplification Shader ← 粗粒度剔除/LOD
     ↓                              ↓            
Vertex Shader                  Mesh Shader        ← 线程组协作处理 meshlet
     ↓                              ↓            
(Hull/Domain/GS)               Rasterizer        
     ↓                       
 Rasterizer    

核心概念	一句话理解
Meshlet	将大网格切成小块，每块就是一个线程组的"工作包"
网格着色器	用计算着色器的方式做几何处理，绕过输入装配器
放大着色器	网格着色器的"守门人"，提前剔掉不可见 meshlet
NGG 快速启动	硬件为网格着色器跳过固定功能的顶点重用和子组构建
GPU 工作图	GPU 自己决定下一步该跑什么着色器，不用等 CPU 下令

硬件加速光线追踪：BVH 构建、降精度遍历与短栈重启机制

---

一、光线追踪概述

1.1 两种可见性求解范式

方法	核心思路	典型应用
光栅化（Rasterization）	遍历每个物体，确定未被遮挡部分	实时渲染主流管线
光线投射 / 光线追踪（Ray Casting / Ray Tracing）	遍历每个像素，向场景投射光线求交	离线渲染、实时光追

• 光栅化天然契合 GPU 的 空间相干性 假设，缓存命中率高
• 光线追踪（尤其是漫反射产生的次生光线）具有 高度发散性与非相干性 ，会严重破坏缓存局部性
• 核心瓶颈从"算力"转移到 显存带宽 和 片上存储（寄存器文件 / LDS）

1.2 全局光照的动机

• 光栅化 + 屏幕空间算法（如 SSR）只能 近似 间接光照，存在信息缺失
• 光线追踪从物理原理出发，天然支持 精确反射、折射、全局光照
• 现代图形 API（DXR / Vulkan RT）将光追提升为与光栅化、计算 同等地位 的核心管线

1.3 DXR 光追管线的可编程阶段

着色器	职责
光线生成着色器（Ray Generation）	调用 TraceRay() 发射光线
相交着色器（Intersection）	自定义几何体求交（非三角形）
任意命中着色器（Any Hit）	处理半透明等中间命中
最近命中着色器（Closest Hit）	表面着色（等价于光栅化中的像素着色器角色）
未命中着色器（Miss）	光线逃逸场景时的背景 / 天空着色

• 固定功能 负责遍历加速结构 + 三角形求交测试
• 关键设计 ：每条光线 独立处理 ，彼此无依赖 → 天然适合 SIMT 并行

---

二、BVH 构建

2.1 加速结构基础

• 包围体层次结构（BVH） ：树形结构，父节点的 AABB 包裹所有子节点
• 遍历时利用空间相干性 逐层剔除 ：光线不与父包围盒相交 → 跳过整棵子树
• 实时场景要求 快速构建 （动态物体每帧重建）+ 高质量树结构 （减少冗余求交）

2.2 LBVH：极速构建

核心思想 ：把三维空间问题降维为一维排序问题

1. 空间映射 ：图元质心 → 莫顿码（Morton Code），交错 XYZ 各位 → 映射到 Z 阶曲线
2. 并行基数排序 ：GPU 高度优化的 $O(N)$ 排序，空间相邻图元在内存中聚拢
3. 层级生成 ：相邻莫顿码的 最高不同位（Highest Differing Bit） 直接决定分割平面，每个线程独立处理一个节点，无需全局同步

优点 ：构建速度极快，接近 $O(N)$，显存访问模式高度可预测

缺点 ：仅依据质心位置聚类，不考虑几何体实际形状和尺寸

• 狭长三角形被强行合并 → 父包围盒大量空白 → 遍历时冗余求交增多

2.3 H-PLOC：质量与速度的平衡

核心思想 ：在 LBVH 框架内嵌入 基于表面积启发式（SAH）的局部聚类

• 局部邻域搜索 ：搜索半径 $R$（如 16 个邻居），并行计算合并后的 SAH 代价
• 双向确认合并（Mutual Nearest Neighbor） ：A 认为 B 最优 且 B 认为 A 最优 → 才合并
• 迭代收敛 ：每层级重复上述过程，直到归约为根节点
• H-PLOC 关键改进 ：将 PLOC++ 的合并逻辑"微缩"塞进 LBVH 的单次自底向上遍历中 → 全程只需 1 次 Kernel Launch

2.4 SBVH：极致质量

• 传统 SAH 的"对象划分"规则：一个图元只属于一个叶节点
• SBVH 引入空间分割 ：允许将跨越空间的三角形 切割成碎片 ，分别放入不同叶节点
• 代价 ：内存增大（节点数增加）、构建极慢（需裁剪几何体）
• 收益 ：最紧凑包围盒、最小空间重叠 → 极致遍历性能
• 适用于 静态场景预计算

2.5 选择策略总结

场景特征	推荐算法	理由
海量微小、均匀分布的图元（粒子）	LBVH	质量劣势被掩盖，极速构建
每帧变形的动态几何	H-PLOC	遍历质量显著优于 LBVH，构建开销仅略高
静态环境	SBVH	一次性构建，运行时最大光追帧率

---

三、降低精度的水密遍历

3.1 问题背景

• 全精度 32-bit float 存储包围盒 → 带宽爆炸 + 芯片面积巨大
• 目标：用 量化压缩 + 低精度计算 节省带宽，但必须解决数值误差引发的 漏检（false misses）

3.2 相对编码与量化压缩

核心思想 ：不存绝对坐标，存相对于父节点的量化偏移

• 父节点包围盒 $[{\mathbf{u}}^{parent},{\mathbf{v}}^{parent}]$
• 局部网格原点对齐 ${\mathbf{u}}^{parent}$，网格尺寸 = 大于父包围盒尺寸的最小 2 幂次 $2^{e_i^{parent}}$
• 量化公式（沿轴 $i$）：

$r_i=\lfloor \frac{(p_i-u_i^{parent})}{2^{e_i^{parent}}}\cdot 2^{N_b}\rfloor ,s_i=\lfloor \frac{(v_i^{parent}-q_i)}{2^{e_i^{parent}}}\cdot 2^{N_b}\rfloor$

• $N_b$ = 量化位宽（如 8-bit），$(\mathbf{r},\mathbf{s})\in [0,2^{N_b}-1]$
• 解压缩 在遍历时迭代进行（从根节点包围盒开始）：

$u_i=u_i^{parent}+r_i\cdot 2^{(e_i^{parent}-N_b)},v_i=v_i^{parent}-s_i\cdot 2^{(e_i^{parent}-N_b)}$

• 层级依赖的压缩 → 极其紧凑的存储

3.3 水密性保证（Watertightness）

核心原则 ：所有误差必须是 保守的 ——包围盒只能变大（误报），绝不能变小（漏报）

具体手段—— 控制浮点舍入方向 ：

计算目标	舍入方向	效果
$t_{min,i}$（光线进入点）	向下舍入 $\mathrm{RD}(\cdot )$	进入点不晚于真实值
$t_{max,i}$（光线离开点）	向上舍入 $\mathrm{RU}(\cdot )$	离开点不早于真实值

• 低精度递推公式（假设 $d_i>0$）：

$t_{min,i}=\mathrm{RD}\left(t_{min,i}^{parent}+\mathrm{RD}\left(\frac{1}{d_i}\cdot r_i\right)\cdot 2^{(e_i^{parent}-N_b)}\right)$

$t_{max,i}=\mathrm{RU}\left(t_{max,i}^{parent}-\mathrm{RD}\left(\frac{1}{d_i}\cdot s_i\right)\cdot 2^{(e_i^{parent}-N_b)}\right)$

• 相交判定加入 ulp 补偿 ：

$\max (t_{min,i})\le \min \left(t_{max,i}+2\mathrm{ulp}(t_{max,i})\right)$

• $\mathrm{ulp}(\cdot )$ = 末位单位（unit in the last place），补偿浮点精度的最后一点误差

总结 ：以微小的包围盒膨胀（保守误报） 换取 零漏检 ，同时用低位宽量化大幅压缩带宽

---

四、多叉树短栈遍历与重启轨迹

4.1 问题背景

• 传统 BVH 遍历需要 完整堆栈 记录"命中但未遍历的子分支"
• 海量并发光线 × 深度 BVH → 片上存储 + 显存带宽 开销巨大
• kd-tree 可用"截断光线 + 从根重启"替代栈，但 BVH 节点空间重叠 → 单纯截断会死循环

4.2 短栈（Short Stack）

• 仅使用 极小容量的栈 （如 4 个条目）
• 绝大多数情况 ：局部遍历几次进出栈就够用，性能极佳
• 特殊情况处理 ：
  • 栈满：丢弃栈底最老条目
  • 栈空但仍有未遍历分支：触发从根节点重启

4.3 重启轨迹（Restart Trail）

核心数据结构 ：为 BVH 每一层分配 1 bit ，记录遍历状态

bit 值	含义
0	近子节点已处理，远子节点 尚未 彻底遍历
1	该层级所有子节点 已处理完毕

工作流程 ：

1. 正常向下遍历 ：选择近子节点处理，远子节点压入短栈，对应层级 bit = 0
2. 出栈 ：自下而上扫描轨迹，找到最近的 0 → 翻转为 1，其下所有层级清零
3. 重启（栈空时） ：从根节点重新向下遍历，直接读取轨迹 而非重算包围盒求交
   • 遇到 bit = 0 → 说明远子节点未处理 → 光线被引导到那个分支

关键优势 ：

• 重启后能 精准空降 到丢失的远子节点，不会死循环
• 上层节点通常已在 L1 Cache 中，且不需重做求交数学 → 额外开销极小
• 寄存器和显存带宽节省 却非常显著

4.4 从二叉树到多叉树（Wide BVH）

• 为适应 GPU SIMD 架构，BVH 常转为 4-wide 或 8-wide 多叉树
• 优势：减少树深度 、同时测试多个包围盒、隐藏内存延迟

多叉树的轨迹升级 ：每层从 1 bit → 计数器（0 到 N）

• counter = k：已处理 k 个相交子节点
• counter = N：当前节点所有子节点处理完毕

遍历逻辑 ：

1. 光线与 N 个子包围盒求交 → 按距离排序得到相交列表
2. 读取计数器 k → 跳过前 k 个已处理节点
3. 选最近的进入遍历，其余压入短栈，更新计数器

4.5 栈剔除（Stack Culling）

场景 ：入栈时命中距离 = 10，出栈时光线已在其他分支找到距离 = 5 的交点 → 该节点无需再遍历

传统方案的困境 ：需要在栈中额外存储每个节点的命中距离 → 加大内存负担

短栈 + 重启轨迹的巧妙方案 ：

1. 当短栈容量紧张时，将 父节点 （而非子节点）压入栈中（用 1-bit 标记）
2. 弹出父节点时，光线的最大命中距离可能已缩短
3. 重新与父节点的 N 个子包围盒求交 → 原先较远的子节点可能直接被剔除

代价 vs 收益 ：

额外代价	收益
1-bit 标记 + 弹出时重做父节点求交	大量冗余三角形求交被剔除
极小的计算开销	显著的带宽与存储节省

4.6 为何完整栈无法简单替代

即使完整栈用 {父节点指针 + 已遍历索引} 替代 {子节点指针 + 距离}，栈深度仍然可能非常大 → 片上存储和访存带宽开销未根本解决。短栈 + 重启轨迹以 极小的重新计算代价 换来了 内存、带宽、剔除率的三重改善 。

---

五、Direct3D 12 渲染管线编程要点

5.1 命令生成与执行分离

传统 API（OpenGL）	现代 API（D3D12）
即时上下文，命令立即提交执行	命令列表（Command List） 先录制，再统一提交
CPU 等待 GPU 完成	命令队列（Command Queue） 异步执行
驱动自动同步	开发者用 围栏（Fence） 手动同步 CPU-GPU
驱动自动管理资源状态	开发者用 资源屏障（Resource Barrier） 手动转换

• 多个 CPU 线程可各自维护命令列表 → 充分利用多核
• 多个命令队列（图形 / 计算 / 复制）→ 充分利用 GPU 异构引擎

5.2 资源与描述符体系

资源创建三种方式 ：

类型	特点
已提交资源（Committed）	类似传统 API，隐式堆
预留资源（Reserved）	仅分配虚拟地址，按需映射物理内存
放置资源（Placed）	预先创建堆，手动放置 → 避免频繁创建/销毁堆

描述符（Descriptor） ≈ 资源视图，9 种类型（CBV / UAV / SRV / Sampler / RTV / DSV / IBV / VBV / SOV）

绑定层次 ：

• 根签名（Root Signature） → 定义着色器期望的数据布局
  • 根常量：直接传值
  • 根描述符：免查表的少量描述符
  • 描述符表：索引到描述符堆

5.3 流水线状态对象（PSO）

• D3D11：多个粗粒度状态对象分别设置 → 每次设置驱动重调硬件
• D3D12 PSO ：统一整合着色器字节码、输入布局、混合 / 光栅 / 深度模板状态、渲染目标格式、根签名等 → 一次性切换

5.4 光追管线的编程差异

维度	光栅管线	光追管线
输出目标	RTV（交换链后台缓冲区）	UAV （需手动复制到后台缓冲区）
根签名	单一根签名	全局根签名 + 局部根签名
加速结构	无	需用顶点 / 索引缓冲区构建 BVH 加速结构
着色器调度	固定阶段顺序触发	相交测试 动态触发 ，配合 着色器绑定表（SBT）
编程模型本质	阶段化管线	更接近 计算流水线

全局根签名 ：所有着色器可见，含渲染结果 UAV + 加速结构 SRV

局部根签名 ：绑定到具体的 hit/miss 着色器，携带材质等信息，与 着色器绑定表 配合

着色器绑定表（Shader Binding Table） ：决定不同光线-几何交互事件应调用哪个着色函数

---

六、总结与展望

硬件光追的核心设计哲学： 以计算换带宽和存储

LBVH / H-PLOC  →  快速构建加速结构
    ↓
宽 BVH（4/8-wide） →  提升 SIMD 并行度，减少树深度
    ↓
相对编码 + 降精度 + 保守舍入  →  压缩带宽，保证水密性
    ↓
短栈 + 重启轨迹 + 栈剔除  →  极低片上状态开销

现实挑战 ：甜品级显卡光追算力仍然有限，纯光栅化性能代际提升放缓

未来趋势 ：光栅化保底 + 稀疏光追 + 张量核心驱动的深度学习重建（DLSS / 帧生成 / 神经渲染）→ 深度异构融合管线