Ch07: GPU 计算入门——从渲染管线到通用计算

Part 2: 移动端推理基础（从零开始） 前置章节：Ch06: ARM 架构与 SIMD 入门 后续章节：Ch08: 内存管理入门

GPU 的前世今生：从画画工具到算数高手

GPU 最初被发明出来只干一件事：画画。更准确地说，是渲染 3D 图形——把游戏里的三角形、纹理、光影效果变成屏幕上的像素。这个任务的特点是：屏幕上有几百万个像素，每个像素的颜色计算逻辑基本相同（采样纹理、计算光照、混合颜色），而且像素之间互不影响——你算你的红色、我算我的蓝色，互不干扰。

为了高效完成这种”大量独立且相似的计算”，GPU 被设计成了一种和 CPU 截然不同的架构：与其像 CPU 那样拥有少量功能强大的核心来处理复杂的逻辑分支，GPU 选择拥有大量功能简单的核心来并行处理简单的计算。

后来人们发现，深度学习的核心运算——矩阵乘法和卷积——恰好也是”大量独立且相似的计算”。一个矩阵乘法 C = A × B 中，C 的每个元素都可以独立计算（互不干扰），且计算逻辑完全相同（取一行乘一列再求和）。于是 GPU 从”画画工具”变成了”算数高手”——这就是 GPGPU（General-Purpose computing on GPU，通用 GPU 计算）的由来。

用一个更直观的类比：CPU 像 4-8 个精英工程师，每个人都能处理复杂的逻辑判断和决策，但总人数少。GPU 像上千个普通工人，每个人只会做简单的重复性工作，但人多力量大——分配任务后，上千人同时开工，总产出远超几个精英。深度学习推理恰好就是”简单重复性工作极多”的场景。

GPU vs CPU：架构的本质区别

为了更清楚地理解 GPU 在推理中的价值，需要在架构层面理解 CPU 和 GPU 的设计取舍。

CPU：为复杂逻辑而生

CPU 的核心设计目标是低延迟——单个任务尽可能快地完成。为此 CPU 拥有：

• 复杂的分支预测器（猜测 if/else 走哪条路，提前开始计算）
• 深层指令流水线（把一条指令拆成多个阶段并行执行）
• 大容量缓存（L1/L2/L3，减少访问内存的等待时间）
• 乱序执行引擎（自动重排指令顺序以消除等待）

这些”聪明”的硬件占了 CPU 芯片面积的大部分。真正用于计算的 ALU（算术逻辑单元）可能只占芯片面积的 10-20%。这就是为什么 CPU 核心数量很少（手机上 4-8 个）——每个核心的”智力成本”太高了。

GPU：为大规模并行而生

GPU 反过来：把”聪明”的硬件压缩到最少，把芯片面积尽量多地分配给计算单元。GPU 核心（在 NVIDIA 术语中叫 CUDA Core，在 ARM Mali 中叫 Shader Core）的控制逻辑非常简单——它们不擅长处理复杂的分支判断，也没有深层乱序执行引擎。但是，一个 GPU 可以有几百甚至上千个这样的核心。

移动 GPU 的规模比桌面 GPU 小得多，但仍然比 CPU 有显著的并行优势。比如高通 Adreno 740（骁龙 8 Gen3）大约有数百个 shader 核心，Apple A17 Pro 的 GPU 有 6 个核心但每个核心内部有大量的 ALU。和手机 CPU 的 4-8 个核心相比，GPU 的”工人”数量依然多出一两个数量级。

什么时候该用 GPU

判断一个计算任务是否适合 GPU 加速，核心标准是两个。

第一，数据并行度够不够高。如果你要处理的数据量很小（比如一个 1x1x1x1 的张量），GPU 的上千个核心大部分在空转——启动 GPU 的开销可能比计算本身还大。反过来，如果你要对一个 128x112x112 的特征图做逐元素操作，有 160 万个元素可以并行处理，GPU 就能大显身手。

第二，计算和数据搬运的比值够不够高。GPU 计算虽快，但把数据从内存搬到 GPU 计算单元也有开销。如果一个操作读取 100 个数据但只做 1 次加法，那搬运的代价远大于计算的收益——这种”低计算密度”的操作在 GPU 上反而不如 CPU。卷积和矩阵乘法的计算密度很高（每读一个权重值，要对多个输入位置做计算），所以非常适合 GPU。逐元素的 ReLU（读一个、比较一次、写一个）计算密度极低，在 GPU 上加速效果有限。

移动 GPU 的三个特殊之处

上一章已经提到移动 GPU 和桌面 GPU 有本质差异。这里从 GPGPU（通用计算）的角度展开。

特殊一：统一内存架构（UMA）

桌面 GPU（比如 NVIDIA RTX 4090）有独立的显存（GDDR6X，24GB），通过 PCIe 总线和 CPU 的系统内存连接。数据要先从系统内存拷贝到显存，GPU 才能访问。这个拷贝操作有延迟也有带宽限制。

移动端完全不同。手机里 CPU、GPU、NPU 共享同一块物理内存（LPDDR5X）。理论上 CPU 写入内存的数据，GPU 可以直接读——不需要跨总线拷贝。

对推理引擎的影响：在移动端，CPU 和 GPU 之间的数据传输成本低得多。这意味着”一些层在 CPU 跑、一些层在 GPU 跑”这种混合执行策略的开销比桌面端小。但代价是 CPU 和 GPU 共享内存带宽——GPU 读数据多了，CPU 可用的带宽就少了。

实际上 UMA 并不意味着完全零拷贝。GPU 通常有自己的缓存体系和地址空间映射。在 OpenCL 和 Vulkan 中，你仍然需要显式管理”主机内存”和”设备内存”之间的映射关系。只是在物理层面，它们指向同一块 DRAM。

特殊二：Tile-Based Rendering (TBR)

几乎所有移动 GPU 都采用 TBR 架构（高通 Adreno、ARM Mali、Apple GPU、Imagination PowerVR 都是）。TBR 的核心思想是把渲染目标（比如屏幕）划分成小块（tile，通常 16x16 或 32x32 像素），每个小块的数据在 GPU 内部的高速缓存（on-chip memory）中处理完毕后，一次性写回主内存。

TBR 对通用计算的影响比较微妙。一方面，如果你的计算访问模式天然具有二维局部性（比如图像处理——相邻像素通常一起被访问），TBR 架构下 on-chip memory 能提供很高的带宽。MACE 大量使用 OpenCL 的 image2d 数据类型（纹理对象），就是为了利用 TBR 的二维缓存优势。另一方面，如果你的计算是全局性的（比如全局平均池化，需要读取整个特征图），TBR 的分块处理反而增加了复杂性。

特殊三：驱动碎片化

这是移动端 GPU 计算最痛苦的现实问题。桌面端基本上只有 NVIDIA（CUDA 生态完善、驱动稳定）和 AMD（ROCm 在追赶）。移动端有四大 GPU 厂商，每家的驱动质量和 API 支持都不同。

同一段 OpenCL 代码在高通 Adreno 上运行正常、在某些 Mali GPU 上可能崩溃或产生错误结果——这在实际开发中是真实发生的问题。驱动 bug、不完整的 API 实现、不一致的性能表现，让移动端 GPU 计算的可靠性远不如桌面端。

这就是 ncnn 选择 Vulkan 而非 OpenCL 的重要原因之一。Vulkan 作为更新的标准，驱动实现的一致性好于 OpenCL。而且 Vulkan 是跨平台的——同一套 Compute Shader 代码可以在 Android、Windows、Linux、macOS 上运行，不需要为每个平台写不同的代码。

GPU 通用计算的三种 API

在移动端做 GPU 通用计算，有三种主要的 API 可选。

OpenCL：最老牌，但碎片化严重

OpenCL（Open Computing Language）是 Khronos Group 2008 年发布的异构计算标准。它的设计目标是统一 CPU/GPU/DSP 等不同计算设备的编程接口。

在移动端，OpenCL 是目前使用最广泛的 GPU 计算 API。MNN、TNN、MACE、Paddle-Lite 都支持 OpenCL 后端。它的编程模型比较直观：你写一段叫做 kernel 的计算代码（用 OpenCL C 语言，类似 C99），指定要启动多少个工作项（work-item），GPU 自动把这些工作项分配到各个核心并行执行。

// OpenCL kernel 示例：两个向量相加
__kernel void vector_add(
    __global const float* a,   // 输入向量 a
    __global const float* b,   // 输入向量 b
    __global float* c,         // 输出向量 c
    const int n)               // 向量长度
{
    int i = get_global_id(0);  // 获取当前工作项的全局ID
    if (i < n) {
        c[i] = a[i] + b[i];   // 每个工作项处理一个元素
    }
}

OpenCL 的问题是”标准很好但驱动很差”。不同厂商的 OpenCL 驱动对标准的实现程度不同，有些功能声称支持但实际使用时有 bug。MNN 的 OpenCL 后端代码中有大量针对特定 GPU 型号的 workaround（绕过 bug 的补丁代码），这些代码的维护成本很高。

Metal：Apple 独占，体验最好

Metal 是 Apple 2014 年推出的图形和计算 API，仅限 iOS 和 macOS 平台。在 Apple 设备上，Metal 是唯一的高性能 GPU 计算路径（Apple 不支持 OpenCL，iOS 上也不支持 Vulkan）。

Metal 的优势是 Apple 完全控制了硬件和驱动——没有碎片化问题。Metal 的 Compute Shader 用 MSL（Metal Shading Language）编写，和 OpenCL C 类似但语法更现代：

// Metal compute shader 示例：两个向量相加
kernel void vector_add(
    device const float* a [[buffer(0)]],  // 绑定到 buffer slot 0
    device const float* b [[buffer(1)]],
    device float* c [[buffer(2)]],
    uint i [[thread_position_in_grid]])    // 当前线程在网格中的位置
{
    c[i] = a[i] + b[i];
}

MNN 和 TNN 都有 Metal 后端用于 iOS 推理。在 Apple 设备上，Metal 的性能通常优于 OpenCL（即使 OpenCL 可用的平台上）。

Vulkan Compute：最新标准，ncnn 的唯一选择

Vulkan 是 Khronos Group 2016 年发布的新一代图形和计算 API，设计目标是解决 OpenGL/OpenCL 的历史包袱，提供更底层、更高效的硬件控制。

ncnn 是六个主流推理引擎中唯一选择 Vulkan 作为 GPU 后端的。这个选择有几个重要原因。

第一，跨平台覆盖。Vulkan 支持 Android、Windows、Linux、macOS（通过 MoltenVK 转译到 Metal），一套 Compute Shader 代码覆盖所有主流平台。相比之下，OpenCL 在 iOS 上不可用、在部分 Android 设备上有 bug；Metal 只在 Apple 设备上可用。

第二，显式控制。Vulkan 把资源管理和同步的控制权交给了开发者（不像 OpenCL 由驱动隐式管理），减少了驱动的”黑盒”行为，bug 更少、性能更可预测。

第三，驱动质量。Vulkan 的 conformance test suite（一致性测试套件）比 OpenCL 更严格，厂商的 Vulkan 驱动实现一致性普遍好于 OpenCL。

代价是 Vulkan 的编程复杂度显著高于 OpenCL——你需要手动管理命令缓冲区、描述符集、管线布局等大量底层概念。ncnn 在内部封装了 VkCompute（命令录制）、Pipeline（shader dispatch 封装）等抽象来简化开发，但编写新的 Vulkan 算子仍然比 OpenCL/Metal 门槛高。

Compute Shader 概念简介

无论用 OpenCL、Metal 还是 Vulkan，GPU 通用计算的编程模型都是类似的。核心概念有两个：Compute Shader（计算着色器）和执行模型。

Compute Shader

Compute Shader 是一段运行在 GPU 上的程序，描述了”每个线程做什么”。它和 CPU 程序的最大区别是：CPU 程序通常描述”按顺序做 A，然后做 B”，Compute Shader 描述”每个工人拿到自己的编号后该做什么”——上千个工人同时执行同一段 Shader 代码，但各自处理不同的数据。

执行模型：Workgroup 和 Thread

GPU 把要执行的工作组织成一个”网格”（Grid），网格由多个”工作组”（Workgroup）组成，每个工作组包含多个”线程”（Thread/Work-item）。

Grid (全局网格)
├── Workgroup 0 [Thread 0, Thread 1, ..., Thread 63]
├── Workgroup 1 [Thread 0, Thread 1, ..., Thread 63]
├── Workgroup 2 [Thread 0, Thread 1, ..., Thread 63]
└── ...

同一个 Workgroup 内的线程可以共享一块高速的本地内存（shared/local memory），还可以做同步（barrier）。不同 Workgroup 之间的线程不能直接通信。这个设计反映了 GPU 硬件的物理结构——一个 Workgroup 映射到一个 GPU 计算单元，本地内存就是该计算单元的片上缓存。

Vulkan Compute Shader 示例

下面是一个用 GLSL（OpenGL Shading Language）编写的 Vulkan Compute Shader，实现两个向量相加：

#version 450                          // GLSL 版本号
layout(local_size_x = 64) in;         // 每个 Workgroup 有 64 个线程

layout(binding = 0) readonly buffer BufA { float a[]; };  // 输入 a
layout(binding = 1) readonly buffer BufB { float b[]; };  // 输入 b
layout(binding = 2) writeonly buffer BufC { float c[]; }; // 输出 c

layout(push_constant) uniform Params {
    int n;                             // 向量长度
};

void main() {
    // gl_GlobalInvocationID.x = WorkgroupID * local_size_x + LocalInvocationID
    uint i = gl_GlobalInvocationID.x;  // 当前线程的全局编号
    if (i < n) {
        c[i] = a[i] + b[i];           // 每个线程处理一个元素
    }
}

这段 GLSL 代码会被编译成 SPIR-V（Standard Portable Intermediate Representation）中间字节码，然后嵌入到应用程序中。运行时 Vulkan 驱动把 SPIR-V 编译成目标 GPU 的机器码执行。

ncnn 的每个 Vulkan 算子都有多个 .comp（compute shader）文件，按 packing（打包数）、算法（direct/Winograd）、数据类型（fp32/fp16）组合成不同变体。编译时这些 .comp 文件被编译为 SPIR-V 嵌入到 ncnn 的二进制库中。

Workgroup 大小的选择

local_size_x = 64 这个值不是随意选的。太小（比如 1）会导致 GPU 调度开销占比过大；太大（比如 1024）可能超过 GPU 硬件的限制或导致寄存器溢出。常见的选择是 64、128 或 256。不同 GPU 硬件有不同的最优值——这又是移动端碎片化问题的一个体现。

ncnn 的 Vulkan 后端通过运行时查询 GPU 的 maxComputeWorkGroupSize 来确定最优的 Workgroup 大小，并在 Shader 中通过 specialization constant（特化常量）传入。

移动 GPU 做推理的优势和劣势

了解了 GPU 计算的基本原理后，来总结移动 GPU 用于推理的利弊。

优势

并行度高，适合大张量运算。 卷积和矩阵乘法的输出元素可以独立计算，天然适合 GPU 的大规模并行。对于计算量大的模型（如 ResNet-50、大尺寸 YOLO），GPU 推理通常比 CPU 快 2-5 倍。

可以释放 CPU 给其他任务。 推理在 GPU 上执行时，CPU 可以同时做图像预处理、用户界面渲染等其他工作。在实时视频处理场景中，这种 CPU/GPU 并行是提升整体帧率的关键。

fp16 加速几乎免费。 移动 GPU 普遍对 fp16 有硬件级支持（fp16 的吞吐量是 fp32 的两倍），而且大多数推理任务在 fp16 精度下不会有明显的精度损失。这意味着在 GPU 上使用 fp16 可以获得接近 2 倍的加速，几乎没有副作用。

劣势

数据搬运开销。 虽然移动端是 UMA（统一内存），但 GPU 计算前仍需要做数据格式转换（比如 NCHW → image2d）和缓存同步。对于小张量，这些开销可能超过计算本身节省的时间。

小张量不划算。 GPU 的优势在大规模并行。如果张量很小（比如 1x32x1x1 的全连接层输出），GPU 的大部分核心在空转，启动 GPU kernel 的固定开销反而让总时间更长。推理引擎通常会对模型中的每一层判断”该在 CPU 还是 GPU 上执行”——这就是 ORT 的 Execution Provider 图分区和 MACE 的 Flow 分流要解决的问题。

功耗通常高于 CPU。 GPU 为了维持大量核心的运行，功耗通常比 CPU（使用效率核时）要高。如果你的应用对功耗敏感（比如后台持续运行的 AI 功能），GPU 可能不是最佳选择——NPU 的能效比通常比 GPU 更好。

编程和调试复杂度高。 编写 GPU Compute Shader 的难度显著高于 CPU 代码。调试工具也不如 CPU 侧成熟——你不能在 Shader 中设断点单步调试（至少在大多数移动端工具中是这样）。这意味着 GPU 后端的开发和维护成本更高。

纹理缓存优化：MACE 为什么这么做

在讨论移动 GPU 时，有一个值得单独讲解的优化技巧：利用 GPU 的纹理缓存（Texture Cache）来加速推理。MACE 是六个引擎中在这个方向上走得最远的。

GPU 的纹理缓存是专门为图像渲染优化的硬件——它的设计假设是”访问模式具有二维空间局部性”（即当你读了像素 (x,y) 后，很可能接下来读 (x+1,y) 或 (x,y+1)）。纹理缓存的命中率在这种访问模式下非常高。

OpenCL 提供了两种数据类型：buffer（一维数组，类似 CPU 内存）和 image2d（二维纹理对象，利用纹理缓存）。如果把推理中的张量映射成 image2d，卷积操作的访问模式（滑窗扫描特征图）恰好匹配纹理缓存的二维局部性假设，可以获得更高的缓存命中率。

MACE 为此定义了 8 种不同的 tensor → image2d 映射方式：

BufferContentType:
  CONV2D_FILTER     → 卷积核的最优 image2d 映射
  IN_OUT_CHANNEL    → 输入/输出特征图的映射
  IN_OUT_HEIGHT     → 按高度维度切分的映射
  IN_OUT_WIDTH      → 按宽度维度切分的映射
  WINOGRAD_FILTER   → Winograd 变换后的映射
  ARGUMENT          → 标量参数的映射
  DW_CONV2D_FILTER  → Depthwise 卷积核的映射
  WEIGHT_HEIGHT     → 全连接层权重的映射

每种算子选择最适合自己访问模式的映射方式。比如卷积核用 CONV2D_FILTER 映射，把输出通道和输入通道分别映射到 image2d 的 x 和 y 维度，这样卷积计算中连续访问不同输出通道的同一个权重值时，恰好是 image2d 的 x 方向扫描——纹理缓存命中率最高。

这种精细的映射策略让 MACE 在某些 GPU 上（特别是 Adreno）的性能优于使用 buffer 模式的 MNN/TNN。但代价也很明显：8 种映射方式意味着大量的格式转换代码，每种 Op 都需要精心选择和适配映射方式，代码维护成本很高。而且不同 GPU 厂商的纹理缓存行为不完全一致——在 Adreno 上最优的映射，在 Mali 上可能并不是最优的。

六个引擎的 GPU 路径选择

将六个引擎的 GPU 策略放在一起对比，可以看到不同的设计哲学：

ncnn 的 Vulkan 路线是最独特的——一套代码覆盖所有平台，但编写 GLSL Compute Shader 的门槛较高。MNN 走的是”全覆盖”路线——每个平台用最原生的 API（Android 用 OpenCL/Vulkan、iOS 用 Metal、服务器用 CUDA），代码量大但每个平台都能达到接近最优。ORT 走的是”委托”路线——自己不写 GPU kernel，通过 Execution Provider 把 GPU 计算交给平台原生的 SDK（CoreML、TensorRT 等），抽象层最厚但扩展性最强。

本章小结

GPU 从渲染管线演变为通用计算平台，其核心优势是大规模并行——上百甚至上千个简单核心同时执行相同的计算。移动 GPU 有三个特殊之处：统一内存架构（CPU/GPU 共享内存）、Tile-Based Rendering（影响缓存行为）、驱动碎片化（同一代码在不同 GPU 上行为不一致）。

移动端 GPU 计算有三种主要 API：OpenCL（最老牌、覆盖最广但碎片化严重）、Metal（Apple 独占、体验最好）、Vulkan（最新标准、跨平台、ncnn 唯一选择）。GPU 计算的编程模型基于 Compute Shader + Workgroup/Thread 层级结构。

GPU 对推理的价值主要体现在大张量的卷积和矩阵乘法上。对于小张量、低计算密度的操作，GPU 可能不如 CPU。推理引擎需要智能地决策每一层应该在 CPU 还是 GPU 上执行。

MACE 的 image2d 纹理缓存优化展示了移动 GPU 的一种深度优化路径——利用 GPU 硬件的纹理缓存特性来加速卷积。这种优化效果显著但增加了大量格式转换的代码复杂度。

下一章将转向另一个移动端推理的核心话题——内存管理。在移动端，”省内存”往往比”跑得快”更重要，因为内存不够你连跑都跑不起来。

导读目录：README.md 上一章：Ch06: ARM 架构与 SIMD 入门——NEON 指令为什么重要 下一章：Ch08: 内存管理入门——为什么”省内存”比”跑得快”更重要