【笔记】GPU 图形与通用计算的整体模型

1. 一句话心智模型

现代 GPU 是面向高吞吐并行工作负载的处理器：图形 API 把它组织成渲染管线，CUDA 把 NVIDIA GPU 组织成线程层级、内存层级和异步任务；二者会使用部分相同的可编程计算资源，也会调用各自需要的专用硬件。

GPU 从图形走向通用计算不是“显卡突然变成另一种硬件”，而是可编程 shader、统一着色器架构和通用计算软件栈逐步把已有并行能力暴露给非图形任务。

2. 先分开五个层次

flowchart TB
    A[业务：游戏、桌面、AI、科学计算] --> B[领域框架：引擎、PyTorch、TensorFlow]
    B --> C[编程接口：OpenGL、Vulkan、Direct3D、CUDA]
    C --> D[驱动、编译器与运行时]
    D --> E[GPU 硬件]
    E --> F[通用计算阵列]
    E --> G[纹理、光栅、ROP、光追、矩阵等专用单元]
    E --> H[显存、缓存与互连]

这几层不能互相替代：

• OpenGL/Vulkan/Direct3D 是图形及部分计算 API，不是 GPU 型号；
• CUDA 是 NVIDIA 的并行计算平台和编程模型，不只是一个函数库；
• Shader 和 CUDA kernel 都是在 GPU 上执行的程序，但接口、执行语义和生态不同；
• PyTorch/TensorFlow 在更高层建模张量与自动微分，通常通过 CUDA 库或其他后端使用 GPU；
• TPU 是针对机器学习工作负载设计的加速器，不是因为 GPU “已经不能叫 GPU”才出现。

3. 图形管线为什么天然适合并行

把三维场景变成屏幕像素，需要对大量顶点、图元和像素执行相似运算：坐标变换、插值、纹理采样、光照、深度测试和混合。这些数据项之间常有较高并行度，促使 GPU 走向宽吞吐架构。

一个简化的现代图形管线：

顶点/索引
  → 顶点着色器
  → 图元装配
  → 光栅化
  → 片元着色器
  → 深度/模板测试与混合
  → 颜色附件/帧缓冲

其中 shader 阶段可编程，而图元装配、光栅化和部分输出操作仍具有固定功能或专用硬件。把整个 GPU 描述成“只会矩阵乘法”会漏掉大量图形专用路径。

4. Texture、Transform、Blend 与 Shader

4.1 Texture 是可采样数据资源

纹理通常是具有格式、维度和采样规则的 GPU 数据资源。它可以装颜色、法线、深度、查找表或一般数据，不等于“贴在模型表面的图片文件”。

纹理单元擅长按坐标读取、过滤和处理边界模式。合成器把窗口 buffer 当作纹理，正是利用这种“按坐标采样并变换”的能力。

4.2 Transform 改变空间关系

顶点和窗口可以通过矩阵完成平移、旋转、缩放与投影。矩阵只是表达变换的工具；GPU 的优势来自对大量独立顶点或像素并行执行这些运算。

4.3 Blend 合并已有结果

Alpha blending 根据源颜色、目标颜色和混合因子组合结果。它常用于透明窗口、阴影和 UI 合成，但还要考虑颜色空间、预乘 alpha 和图层顺序。

4.4 Shader 是运行在可编程阶段的程序

“着色器”来自早期计算光照和颜色的用途，后来扩展为多个可编程阶段。顶点 shader、片元 shader 和 compute shader 面向不同输入输出模型。

Shader 不是“GPU 的全部计算原语”。应用通过 API 创建资源、管线和命令，驱动把 shader 中间表示编译成目标 GPU 指令，GPU 再结合固定功能单元执行完整工作。

5. OpenGL、Vulkan、Direct3D 与 CUDA

5.1 图形 API 的共同目标

OpenGL、Vulkan 和 Direct3D 都允许应用使用 GPU 完成图形工作。DirectX 是微软的一组多媒体 API，讨论 3D 图形时通常实际指 Direct3D。

它们都不是直接向每一种 GPU 写机器指令。API、shader 语言/中间表示、用户态驱动和内核驱动共同形成硬件抽象层。

5.2 OpenGL 与 Vulkan 的关键差异

OpenGL 采用较多隐式状态和驱动管理，应用入口相对直接；Vulkan 把资源、同步、队列、command buffer 和 pipeline 等控制更多交给应用。

Vulkan 并不是“没有抽象、直接操作硬件”，也不能简单称为 OpenGL 的新版本。它们是 Khronos 维护的不同 API，抽象和兼容模型不同。Vulkan 的显式模型有助于降低部分驱动期开销和组织多线程命令生成，但开发复杂度更高，也不保证任何程序都自动更快。

Direct3D 12 与 Vulkan 都属于更显式的一代 API，但平台、shader 工具链、资源模型和生态不同。macOS 原生主要使用 Metal；Vulkan 应用可以借助 MoltenVK 映射到 Metal，但这不是原生 Vulkan 驱动。

5.3 图形 API 也能做计算

现代 OpenGL、Vulkan 和 Direct3D 都提供 compute shader。它绕开传统绘制入口，通过工作组执行通用计算，并能与图形资源直接协作。

因此，“图形 API 只画图、CUDA 只计算”是便于入门的近似，不是严格边界。更准确的区分是：

• 图形 API 的资源、同步和管线模型同时服务渲染与计算；
• CUDA 围绕 NVIDIA GPU 的通用并行计算建立了语言扩展、编译器、runtime、库和调试工具生态。

5.4 CUDA 与图形互操作

CUDA 可以与 OpenGL、Vulkan、Direct3D 等共享或导入图形资源，避免不必要地把大块数据复制回 CPU 内存。但“共享显存对象”仍需要显式所有权、布局和同步处理，不能笼统等同于零成本。

6. 从 Shader 到 GPGPU，再到 CUDA

早期通用计算曾把数据编码成纹理，借图形管线和 shader 完成非图形运算。这种方法证明了 GPU 的并行计算价值，但需要把通用问题伪装成图形问题。

随后出现更直接的 GPGPU 接口：

• 图形 API 增加 compute shader；
• CUDA 为 NVIDIA GPU 提供 C/C++ 语言扩展、编译器、runtime 和高性能库；
• OpenCL、SYCL 等提供不同程度的跨厂商计算抽象；
• 机器学习框架继续在上层提供张量、算子图、自动微分和分布式执行。

所以演进主线不是“shader 被 CUDA 替代”，而是同一硬件获得了多种面向不同领域的编程入口。

7. CUDA 源码怎样变成 GPU 工作

CUDA 源文件可以同时包含 host code 与 device code：

__global__ void add(const float* a, const float* b, float* c) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    c[i] = a[i] + b[i];
}

7.1 编译阶段

nvcc 作为编译驱动拆分和协调主机、设备代码：

• host code 交给宿主 C/C++ 工具链生成 CPU 代码；
• device code 可生成 PTX，也可由 ptxas 生成特定 SM 架构的 cubin；
• fatbin 可以同时携带多个 cubin 和 PTX，以覆盖不同 GPU。

PTX 是面向 NVIDIA 虚拟 GPU ISA 的中间表示，不是 Java 字节码的完全等价物。它服务于 GPU 架构兼容和 JIT，但运行模型、类型系统和 VM 语义都不同。

7.2 加载与 JIT

进程通过 CUDA Runtime API 或 Driver API 初始化上下文、分配内存、加载 module。若 fatbin 中已有适配当前 GPU 的 cubin，驱动可以直接选择；若依赖 PTX，驱动会把它 JIT 编译成目标 GPU 机器码，并可能缓存结果。

因此，不能笼统说“device code 总在启动时从 PTX 编译”，也不能说“编译后永远只有一种显卡机器码”。产物组合由编译目标决定。

7.3 提交执行

host code 通过 kernel launch 描述执行规模：

add<<<gridDim, blockDim, dynamicSharedMemory, stream>>>(a, b, c);

这次调用主要是向指定 stream 提交工作。对 host 而言，kernel launch 通常是异步的；结果是否可用要看后续依赖、事件、内存复制和显式同步。

8. Grid、Block、Thread、Warp 与 SM

CUDA 的软件层级：

Kernel launch
  └─ Grid
      └─ Thread Block
          └─ Thread

硬件执行时还要加入两个关键概念：

• Warp：SM 调度线程的基本分组；当前 NVIDIA CUDA 模型中通常是 32 个线程。
• SM（Streaming Multiprocessor）：容纳并调度多个 warp，包含执行单元、寄存器文件、shared memory 等资源。

不能把 CUDA thread 直接等同于 CPU 线程，也不能把 CUDA Core 等同于可独立运行任意线程的 CPU 核心。一个 thread 是编程模型中的执行实例；warp 是调度与执行的重要粒度；底层指令再映射到 SM 内不同执行管线。

8.1 Block 为什么重要

一个 block 被分配给一个 SM 执行，不会跨多个 SM 拆开。block 内线程可以使用 shared memory 和 block 级同步协作。不同 block 原则上必须能独立执行，调度顺序不应成为正确性前提。

8.2 Warp 分歧

同一 warp 中线程执行不同控制流分支时，硬件需要分别推进分支路径并屏蔽不参与的线程，降低有效吞吐。现代 GPU 的独立线程调度增强了灵活性，但没有消除 warp 内控制流分歧的成本。

8.3 Occupancy 不是越高越好

一个 SM 能同时驻留多少 block/warp，受寄存器、shared memory、线程数和架构上限共同约束。较高 occupancy 有助于在某个 warp 等待内存时切换到其他就绪 warp，但最终性能还取决于访存、指令混合和算法结构。

9. GPU 的“时分复用”是什么

GPU 确实会在多个可运行 warp 之间调度，以隐藏长延迟；多个 block 也会随着 SM 资源释放而分批驻留。它与 CPU 抢占式线程调度有相似目标，但机制不同：

• 大量 warp 的寄存器状态可以同时驻留在 SM；
• warp scheduler 从就绪 warp 中选择指令发射；
• 切换通常不需要像 CPU 线程那样保存和恢复完整软件上下文；
• block 一旦驻留，通常会占用资源直到完成，但现代 GPU 还存在更复杂的抢占能力，不能概括成“绝不抢占”。

Device code 的机器码可以在 context/module 生命周期内保持已加载；某次 kernel launch 创建的 grid、block、thread 执行状态只在该次工作期间存在。二者不能都叫“kernel 常驻”。

10. 多个 Kernel 如何并发

同一 stream 中的操作按 stream 顺序建立执行关系；不同 stream 的工作在依赖和硬件资源允许时可以重叠，包括 kernel 与 kernel、计算与内存传输。

“提交并发”不保证“实际同时执行”。是否重叠取决于：

• GPU 是否支持相应并发能力；
• kernel 是否占满 SM、寄存器或 shared memory；
• stream、事件与同步依赖；
• 内存带宽和复制引擎；
• MPS、MIG、进程上下文和系统调度策略。

把 Kernel/Block/SM 类比成 Job/Pod/Node 可以帮助建立层级感，但不能用于推导调度语义。Kubernetes 做分布式资源编排，GPU 硬件调度器处理的是片上执行资源，两者的故障、抢占和生命周期模型完全不同。

11. CUDA 内存层级

从编程视角看，常见存储包括：

存储	典型范围与特点
Register	每线程使用，容量有限，最快但会影响驻留度
Local memory	线程私有语义，但通常落在设备内存并经缓存
Shared memory	block 内共享、软件管理的片上存储
Global memory	所有线程可访问的设备内存，容量大、延迟高
Constant/texture path	面向特定访问模式的只读或采样路径

性能不只取决于“数据在显存还是内存”。访问是否合并、缓存命中、bank conflict、数据复用和计算访存比同样关键。

Host 与 device 分离的系统通常需要显式或隐式搬运数据。统一虚拟寻址统一的是地址管理；managed memory 提供自动迁移；物理共享内存允许某些零拷贝路径。这些概念都不保证访问成本消失。

12. 统一内存的三个不同语境

“统一内存”至少可能指：

1. 共享物理内存架构：CPU、GPU 使用同一物理内存池，例如许多 SoC 和集成 GPU；
2. 统一虚拟寻址：CPU 与 GPU 地址空间具有统一映射关系；
3. CUDA Unified Memory：由运行时和驱动管理数据放置与迁移的编程模型。

Apple Silicon 的 CPU、GPU 等处理单元共享封装内的高带宽内存池，能减少离散 GPU 场景中的显式 PCIe 往返，并让容量动态服务多个处理单元。但“共享物理内存”不等于任何访问都零复制、零同步或等延迟：缓存一致性、资源布局、API 存储模式和内存带宽仍然存在。

PC 并非完全没有共享内存架构。集成 GPU 和 APU 长期共享系统内存；Apple 的突出点是软硬件统一设计、高带宽封装和产品化取舍。模块化、可升级性、成本、容量和带宽之间存在工程权衡，不应简化为某一家厂商“无法跟进”。

13. Tensor 与 TPU 放在什么位置

机器学习框架中的 tensor 通常表现为带数据类型、shape、stride、device 等元数据的多维数据容器。标量、向量、矩阵可以分别用 0、1、2 维 tensor 表示，但数学中的张量还强调坐标变换规律，不能在所有语境下简单等同于 N 维数组。

TPU 针对神经网络中的大规模张量运算提供专用数据路径和矩阵计算单元，追求的是特定工作负载的吞吐和能效。GPU 仍保留更广泛的图形和通用并行能力；专用加速器的出现是体系结构特化，而不是 GPU 名称失效。

14. AI 框架怎样落到 GPU

PyTorch、TensorFlow 等框架提供张量算子、自动微分、模型组件和设备抽象。以 NVIDIA 路径为例，一次高层算子可能落到：

PyTorch / TensorFlow 算子
  → 框架调度与计算图
  → cuBLAS / cuDNN / 自定义 CUDA kernel
  → CUDA Runtime / Driver
  → GPU

模型 checkpoint、safetensors、GGUF、推理引擎和 GPU 执行是不同层：文件格式描述如何保存权重和元数据，框架解释模型结构，推理引擎组织算子与缓存，CUDA 等后端最终提交 GPU 工作。不能从“文件能装进显存”推导“任意 runtime 都能执行它”。

15. 最容易混淆的结论

• GPU 擅长高吞吐并行，不等于任何矩阵或任何并行任务放上去都会更快。
• Shader 与 CUDA kernel 都可运行在 GPU 计算资源上，但不是同一种 API。
• Compute shader 让图形 API 能做通用计算，图形与计算不是硬边界。
• Vulkan 更显式，不等于无驱动抽象，也不保证天然比 OpenGL 快。
• CUDA thread 是轻量执行实例，不等于 CPU 线程。
• CUDA Core 是执行资源，不等于独立 CPU 核心。
• PTX 类似中间 ISA，但不能直接等同于 JVM 字节码。
• Kernel launch 通常对 host 异步，结果可用性仍需同步关系。
• Device code 可保持加载，执行实例不会永久常驻。
• 统一内存减少部分搬运，不会消灭同步、带宽和数据布局成本。

16. 复习索引

• 两种视角：图形 API 用渲染管线组织 GPU；CUDA 用并行线程和内存模型组织 NVIDIA GPU。
• 演进链：固定图形功能 → 可编程 shader → GPGPU → compute API 与领域框架。
• 编译链：CUDA 源码 → host code + PTX/cubin → fatbin → Driver 选择或 JIT。
• 执行链：launch → grid → block → thread；硬件侧由 SM 调度 warp。
• 并发条件：不同 stream 只提供并发机会，资源和依赖决定是否真实重叠。
• 内存边界：地址统一、自动迁移和物理共享是三个不同概念。
• 框架位置：PyTorch/TensorFlow 在 CUDA 之上建模张量和神经网络。

17. 参考资料

• NVIDIA CUDA Programming Guide
• NVIDIA CUDA Compiler Driver NVCC
• Khronos Vulkan Specification
• Khronos OpenGL 4.6 Specification
• Apple Metal Documentation