AI Compute Stack 全景 — 从推理框架到硬件指令集

为什么这些概念特别绕？

当你听到 CUDA、OpenCL、SYCL、ROCm、oneAPI 这些名字时，最容易犯的错误是把它们当成”一个东西”。实际上，它们每一个都是 “一套东西”的品牌名，横跨编程语言、编译器、运行时、算子库等多个层次。

比如 OpenCL，它既有编程语言（OpenCL C）又有运行时（OpenCL Runtime）；SYCL 是一门语言，但它没有自己的运行时，而是借用别人的后端（OpenCL Runtime / Level Zero / CUDA Runtime）。这种”一个名字覆盖多层”的设计，正是混淆的根源。

本文的目标：用一张 7 层全景图，把每一层是什么、为什么存在、跟谁对话讲清楚。读完之后，你再遇到任何 GPU 计算技术名词，都能精确地把它放到正确的层次上。

Section 0: 全景概览

先建立全局视野。下面这张交互式栈图展示了 AI/GPU 软件栈的 7 个层次，从最上层的推理框架到最底层的硬件指令集。点击任意层可以展开查看该层包含的具体技术；底部品牌按钮可以高亮某个品牌横跨的所有层。

点击层名展开技术节点 · 底部按钮按品牌高亮

一次 MatMul 的旅程

理解各层关系最好的方式，是追踪一次具体操作的完整调用链。下面我们跟踪 model.forward() 中一次矩阵乘法，看它如何从框架层一路下沉到硬件执行：

1. Framework: model.forward() 遇到 MatMul

接下来，我们自底向上逐层展开。

Section 1: Hardware ISA — GPU 能执行的唯一东西

为什么需要这层： 这是 GPU 能直接执行的最底层 — 二进制机器指令。

ISA（Instruction Set Architecture）是 GPU 硬件能理解的指令集，类比 CPU 的 x86 或 ARM 指令集。每个 GPU 厂商的 ISA 互不兼容：

ISA	厂商	特点
SASS	NVIDIA	每代 GPU 架构不同（Ampere、Hopper…），不公开完整文档
RDNA ISA	AMD	RDNA 3、RDNA 4 各有不同
Gen ISA	Intel	Arc / Data Center GPU（Xe 是架构品牌名，ISA 叫 Gen ISA）
Apple GPU ISA	Apple	M1/M2/M3 系列，完全不公开

为什么你不直接写 ISA？ 太底层、厂商不公开完整文档、换硬件就得重写全部代码。

PTX vs SASS（NVIDIA 特有）： PTX（Parallel Thread Execution）是 NVIDIA 设计的虚拟 ISA — 一种稳定的中间层，跨 GPU 代际不变。SASS 是真实硬件 ISA，每代不同。Driver 负责最终的 PTX → SASS 翻译。这让你的 CUDA 程序在新 GPU 上也能跑（前向兼容）。

需要一层帮我们屏蔽硬件差异 → Driver

Section 2: Driver — 不只是”驱动程序”

为什么需要这层： 屏蔽硬件差异，给上层提供稳定接口。

GPU Driver 的角色类似操作系统中的设备驱动，但它做的事比你以为的多得多：

内含编译器后端

这是很多人不知道的 — Driver 不只是”让硬件工作”的驱动，它有完整的编译器后端，负责把中间表示（IR）翻译成硬件 ISA：

厂商	IR → ISA	编译模式
NVIDIA	PTX → SASS	AOT（构建时用 `ptxas` 预编译）或 JIT（运行时由 Driver 编译，支持前向兼容新 GPU）
Intel	SPIR-V → Gen ISA	JIT（Driver 内含 SPIR-V → Gen ISA 编译器）
AMD	LLVM IR / AMD IL → RDNA ISA	主要 AOT（编译时生成目标 ISA），也支持 JIT

硬件资源管理

GPU 显存分配/回收
计算单元调度
多进程 GPU 隔离

Driver 接口太底层（ioctl 级别），需要更友好的编程抽象 → Runtime

Section 3: Runtime — GPU 的 “C Runtime”

为什么需要这层： 给程序员一个”操作 GPU”的编程接口，封装 Driver 的底层细节。

正确理解 Runtime

GPU Runtime 的角色类似 C Runtime (libc) 或 JRE：

C Runtime 给你 malloc/free + pthread_create
GPU Runtime 给你 bufferAlloc/free + kernelDispatch

它不是操作系统（那是 Driver 的角色），而是编程语言级的运行时库。

Runtime 的五大职责

1. Device Discovery — 查询系统有哪些 GPU、它们的能力

2. Context（上下文） — GPU 编程的”会话/沙箱”。Context 是一个容器，绑定了：一个 GPU 设备 + 它名下所有的 buffer、queue、program、kernel。多个进程各自创建独立 Context，资源互不干扰。类比：数据库 Connection — 你在 Connection 上创建 Transaction、执行 Query，Connection 就是 Context。

CUDA Runtime API 隐式管理 Context（你不需要手动创建），但底层仍然存在。OpenCL / Level Zero / Vulkan 都要求显式创建 Context。

3. Buffer 管理 — 在 GPU 显存上分配/释放内存块（类比 CPU 的 malloc/free）

4. Command Queue / Command List：

GPU 不像 CPU 那样”调一个函数就立即执行”，而是通过”命令提交”的方式工作。有两种模型：

Command Queue（CUDA / OpenCL 风格）：即时提交 — 每次 API 调用立即将一条命令放入队列，GPU 按顺序执行。你不能”重放”同一个命令序列，想再跑一遍就得重新调用 API
Command List（Level Zero / Vulkan / Metal 风格）：录制-提交 — 先把一整套命令录制到一个 command list 里（此时 GPU 什么都没做），然后一次性 submit 给 GPU 执行。好处是 driver 可以在录制时预编译和预验证整个命令序列，提交时开销极小。同一个 command list 可以重复提交（重放），适合每帧渲染、迭代训练等场景

Command list 的参数在录制时绑定。如果需要改参数，可以：(1) 录制新的 command list；(2) 用 push constants（Vulkan）在提交时塞少量参数；(3) 用 indirect dispatch — 参数从 GPU buffer 动态读取。

5. Kernel Dispatch — 把编译好的 kernel 程序加载到 GPU，配置线程网格（grid/block），提交执行

Kernel Dispatch 具体流程（OpenCL 为例）

1. Platform/Device discovery — clGetPlatformIDs() → clGetDeviceIDs()
2. Context + Queue — clCreateContext() → clCreateCommandQueue()
3. Buffer — clCreateBuffer(size_A)              // 类比 malloc
4. 数据传输 — clEnqueueWriteBuffer(queue, bufA, hostA)
5. Kernel 加载 — clCreateProgramWithIL(spirv) → clCreateKernel("matmul")
6. 参数绑定 + Dispatch — clSetKernelArg() → clEnqueueNDRangeKernel()
7. 同步 + 读回 — clFinish() → clEnqueueReadBuffer()

这个流程将 Context、Buffer、Queue、Kernel、Dispatch 的关系串联起来。CUDA Runtime API 做了更多隐式封装（自动管理 Context），但底层逻辑相同。

Extension（扩展）

Vulkan / OpenCL 的核心规范定义了基础功能，但 GPU 能力在快速演进。Extension 机制允许在不改动核心规范的情况下增加新能力：

跨厂商标准扩展（VK_KHR_xxx）：Khronos 官方批准，多厂商实现。例如 VK_KHR_ray_tracing_pipeline 给 Vulkan 加了光线追踪
厂商私有扩展（VK_NV_xxx、VK_AMD_xxx）：单厂商实现，可能后续被标准化
使用前必须运行时查询 GPU 是否支持该扩展，支持才能启用

这类似浏览器的 Web API — 核心规范是 baseline，新特性先作为实验性 API 提供，成熟后纳入标准。

主要 Runtime 对比

Runtime	厂商	特点
CUDA Runtime API	NVIDIA	最高层抽象，隐式上下文管理，NVIDIA GPU 开发首选
CUDA Driver API	NVIDIA	更底层，显式控制上下文/模块
OpenCL Runtime	Khronos	跨平台，显式 buffer/queue 管理
Level Zero	Intel	低开销、显式控制，oneAPI 底层 runtime。规范厂商中立，但实际仅 Intel 实现
Vulkan (Compute)	Khronos	极致显式控制，command buffer 录制。不只是图形 API — llama.cpp 就用它做跨平台 GPU 推理
Metal	Apple	Apple 专属，command buffer 模式
HIP Runtime	AMD	ROCm 的 runtime，API 几乎对齐 CUDA Runtime

常见误解

CUDA Runtime API vs CUDA Driver API：同一厂商的两层抽象。Runtime 更简单（隐式 context），Driver API 更底层。大多数人只用 Runtime API
Vulkan 不只是图形 API：Vulkan Compute 可以跑通用计算 kernel
Level Zero vs OpenCL Runtime：都能驱动 Intel GPU，Level Zero 更新更底层（类似 Vulkan 设计哲学）

Section 4: Language → Compiler → IR → Kernel

为什么需要这层： 你需要用某种语言写出 GPU 能跑的程序（kernel），编译器把它翻译成 Runtime 可以 dispatch 的格式。

四个概念

Language — 写 GPU 代码的编程语言/扩展
Kernel — 编译后可被 Runtime dispatch 到 GPU 并行执行的计算程序。它自己不知道怎么跑 — 需要 Runtime 来 allocate buffer、dispatch 它
Compiler — Language → IR 的翻译器（nvcc, DPC++, hipcc, clang…）
IR (Intermediate Representation) — 编译后的中间字节码，类比 Java bytecode — 平台无关，由 Driver 的编译器做最终翻译成硬件 ISA

Single-Source vs Dual-Source

影响开发体验的关键架构差异：

Single-source（CUDA C++, SYCL, HIP）： host 代码和 kernel 写在同一个文件，编译器分别提取。开发体验好，可共享类型定义
Dual-source（OpenCL C, GLSL, HLSL）： kernel 单独写在文件/字符串中，host 代码通过 Runtime API 加载。灵活但体验割裂

HIP 的跨平台机制

HIP 是理解”语言层如何实现跨平台”的好例子：

HIP 语法和 CUDA C++ 几乎一致（hipMalloc ↔ cudaMalloc）
编译器 hipcc 检测目标平台：AMD GPU → HIP-Clang (LLVM) 生成 AMDGCN；NVIDIA GPU → nvcc 生成 PTX
这是源码级可移植性：同一份代码，编译时选择不同后端

GPU 编程语言全景

Language	生态	编译目标 (IR)	特点
CUDA C++	NVIDIA	PTX	NVIDIA 专属，最成熟生态
HIP	AMD (ROCm)	AMD GPU IR / PTX	AMD 对标 CUDA 的语言，语法几乎一致
OpenCL C	Khronos	SPIR-V	跨平台，C99 风格，较老
SYCL	Khronos	SPIR-V (via DPC++)	现代 C++ single-source，Intel 主推
Triton	OpenAI	Triton IR → MLIR → LLVM IR → PTX/AMDGCN	Python 风格写 kernel，自动 tiling
GLSL	Khronos	SPIR-V	图形着色器语言，也可用于 compute shader
HLSL	Microsoft	DXIL / SPIR-V	DirectX 着色器语言
WGSL	W3C	SPIR-V / HLSL / MSL (via Tint/Naga)	WebGPU 着色器语言
Metal SL	Apple	Metal IR (AIR)	Apple 专属
Slang	Khronos 开源	SPIR-V / HLSL / MSL / CUDA / GLSL	新一代跨平台着色器语言，多后端输出

Shader vs Kernel — 两条通道，同一个硬件

这两个词来自不同的历史传统，但最终跑在完全相同的 GPU 硬件（SM/CU）上：

Kernel：通过 Compute API（CUDA、OpenCL、Level Zero、HIP）提交的通用计算程序。这些 API 专为并行计算设计，不涉及图形管线
Shader：通过 Graphics API（Vulkan、Metal、DirectX 12、WebGPU）提交的 GPU 程序。名字来自图形渲染 — GPU 最初只做”着色”（shading），管线中的可编程阶段叫 shader（vertex shader、fragment shader 等）
Compute Shader：Graphics API 中独立于图形管线的计算通道。它功能上等同于 kernel，但通过 Graphics API 的 compute pipeline 提交，不经过 vertex/rasterize/fragment 等图形阶段

为什么有人用 compute shader 而不用 kernel？因为 Vulkan/Metal 的跨平台覆盖面比 CUDA 广得多 — llama.cpp 就用 Vulkan compute shader 实现跨 NVIDIA/AMD/Intel/手机 GPU 的推理。

IR 对比

IR	对应语言	消费者	特点
PTX	CUDA C++	NVIDIA Driver	NVIDIA 专有虚拟 ISA，文本格式可读
SPIR-V	OpenCL C, SYCL, GLSL, HLSL, WGSL, Slang	OpenCL RT, Vulkan, Level Zero	Khronos 标准，二进制格式，跨平台通用
DXIL	HLSL	DirectX 12 Driver	Microsoft 专有
Metal IR (AIR)	Metal SL	Metal Driver	Apple 专有
LLVM IR	Triton, HIP	各厂商 LLVM 后端	通用编译器 IR，被多个工具链复用

Section 5: Operator Library — 你不用手写 Kernel

为什么需要这层： 手写 kernel 太难。算子库提供预优化的 kernel 集合 + 调用 Runtime 的胶水代码。

对上层暴露 matmul(A, B, C) 接口，内部选择最优 kernel、配置 tiling 策略、通过 Runtime API allocate buffer 并 dispatch。矩阵拆分（tiling）就发生在这里 — 大矩阵被切成适合 GPU shared memory 的 tile。

正确类比：算子库 ≈ Intel MKL / BLAS — 性能优化的算法库，不是”标准库”（范围太广了）。

算子库	厂商	Runtime 依赖	覆盖算子
cuDNN	NVIDIA	CUDA Runtime	卷积、归一化、Attention
cuBLAS	NVIDIA	CUDA Runtime	矩阵乘法、BLAS 运算
oneDNN	Intel	OpenCL RT / Level Zero / CPU JIT	卷积、MatMul、归一化
MPS	Apple	Metal	矩阵乘法、卷积、图像处理
XNNPACK	Google	CPU 直接调用	移动端 CPU 优化算子
rocBLAS / MIOpen	AMD	HIP Runtime	BLAS / 深度学习算子

oneDNN 内部怎么工作

oneDNN 是理解”算子库如何使用 kernel + runtime”的最好例子。它支持多个后端（backend），每个后端是一整套针对特定硬件的 kernel 实现 + 对应的 runtime 调用：

Intel GPU 后端：kernel 用 OpenCL C 编写，或用 nGen 在运行时生成 → 通过 OpenCL Runtime 或 Level Zero 提交到 GPU
CPU 后端：使用 Xbyak（x86）/ Xbyak_aarch64（ARM）在运行时生成优化的机器码
NVIDIA GPU 后端：通过 CUDA Runtime 提交 kernel
AMD GPU 后端：通过 HIP Runtime 提交 kernel
选择逻辑：oneDNN 根据输入 tensor 的形状、数据类型、当前硬件，自动选择最优的后端和 kernel 实现

nGen 是 oneDNN 内部的 Intel GPU JIT 汇编器（不是编译器）。它不接受任何编程语言输入，而是一个 C++ API — 开发者用 C++ 函数调用逐条拼出 Intel GPU 指令（类似 CPU 端的 Xbyak 拼 x86 汇编），运行时直接输出二进制。nGen 生成的是 Gen ISA 真实指令（不是 PTX 那样的虚拟 ISA），但 API 屏蔽了不同 Gen 代际之间的编码差异（Gen9 vs Gen12 等），所以同一套 C++ 代码可以为不同代 Intel GPU 生成正确的二进制。

Xbyak 是同样思路的 CPU 版本 — x86/x64 的 JIT 汇编器。oneDNN 在 CPU 端用它根据运行时检测到的 CPU 特性（AVX-512? AVX2?）动态生成最优机器码，比”编译时生成多个版本 + 运行时选择”更灵活。

这展示了算子库的本质：kernel 集合 + runtime 胶水 + 自动选择策略。

Triton 的特殊位置

Triton 介于手写 kernel 和算子库之间 — 你用 Python 风格写 kernel 逻辑，Triton 编译器自动做 tiling 和优化。PyTorch 2.0+ 的 torch.compile 后端大量使用 Triton 生成 kernel。

Section 6: Inference Framework + Graph Optimizer

为什么需要这层： 你不想手动调算子库的 API。推理框架加载模型文件，做图优化，把每个算子 dispatch 到对应的算子库/后端。

ONNX 格式 vs ONNX Runtime（常见混淆）

ONNX — 开放的模型交换格式（.onnx 文件），类比 HTML
ONNX Runtime — Microsoft 的推理引擎，类比 Chrome
其他引擎（TensorRT、OpenVINO）也能消费 .onnx 文件

框架做什么

模型加载 — 解析 .onnx / .tflite / .gguf 等模型文件
图优化 — 算子融合（MatMul+BiasAdd+ReLU → 一个 fused kernel）、常量折叠、layout 转换
调度 — 把图中每个算子 dispatch 到对应的后端

框架	输入格式	后端机制	典型调用链
ONNX Runtime	.onnx	Execution Provider 插件	→ CUDA EP → cuDNN → CUDA RT
TensorRT	.onnx / .plan	NVIDIA 专有引擎	→ 自有 kernel → CUDA RT
OpenVINO	多种→内部 IR	内置插件	→ oneDNN → OpenCL/L0
LiteRT (TFLite)	.tflite	Delegate 插件	→ GPU delegate → OpenCL/Vulkan
CoreML	.mlmodel	Apple 专有	→ MPS / ANE
llama.cpp	.gguf	ggml backends	→ ggml → CUDA/Metal/Vulkan

图优化器 = 图级编译器

TensorRT、XLA、Apache TVM 不是简单的推理框架 — 它们是图级编译器：输入计算图，输出优化后的 kernel 调用序列。做的事：算子融合、内存规划、精度优化（FP16/INT8）、kernel 自动选择。

llama.cpp / ggml 的垂直整合

传统分层（以 NVIDIA 为例）：

TensorRT → cuDNN/cuBLAS → CUDA Runtime → Driver

每层各管各的，标准接口隔开。cuDNN 只服务 CUDA 一个后端；TensorRT 不关心 cuDNN 内部怎么实现 kernel；换 GPU 代际时，只需 cuDNN 更新 kernel，TensorRT 不用改。

ggml 的做法 — 把”算子库”整层吞掉，自己给每个硬件后端从头写 kernel：

llama.cpp → ggml ──┬── ggml-cuda.cu   → CUDA Runtime → Driver
                   ├── ggml-metal.m   → Metal        → Driver
                   ├── ggml-vulkan.cpp → Vulkan       → Driver
                   └── ggml-cpu.c     → (直接执行)

ggml 和 cuDNN/oneDNN/MPS 是同一层的东西 — 都是算子库。区别在于：ggml 定义了一套自己的算子接口（针对 LLM 推理：矩阵乘法、softmax、RoPE、Q4_K_M 量化等），然后为这套接口实现了多个后端。它不用 cuDNN，不是因为层级不同，而是因为 cuDNN 的算子接口不匹配 ggml 的需求（比如 cuDNN 没有 ggml 需要的特殊量化 kernel），加上 ggml 追求零依赖部署。

本质上，ggml 是一个自带跨平台策略的算子库 — 以”为每个后端从头写 kernel”的方式实现跨平台，而不是像 oneDNN 那样复用标准 runtime 的抽象层。

好处：极致控制、零第三方依赖、一个二进制跑所有平台、深度定制（自定义量化格式 GGUF、激进的算子融合）
代价：每加一个硬件后端，所有算子都要从头实现一遍。cuDNN 的 matmul kernel 经过 NVIDIA 工程师多年优化，ggml 的 CUDA kernel 未必能打平

Section 7: 跨层品牌解剖

核心困惑终极解答：这些名字不是”一个东西”，而是”一套东西”的品牌名。

下面用交互式工具，选择不同场景看技术栈路径：

选择一个场景，查看技术栈路径

五大品牌 × 六层对比

层	CUDA (NVIDIA)	ROCm (AMD)	oneAPI (Intel)	OpenCL (Khronos)	Metal (Apple)
Language	CUDA C++	HIP	SYCL (DPC++)	OpenCL C	Metal SL
Compiler	nvcc / NVRTC	hipcc (Clang)	DPC++ / ICX	各厂商实现	Metal Compiler
IR	PTX	LLVM IR→AMDGCN	SPIR-V	SPIR-V	AIR
Runtime	CUDA RT	HIP RT (ROCr)	Level Zero	OpenCL RT	Metal
Operator Lib	cuDNN/cuBLAS	MIOpen/rocBLAS	oneDNN/oneMKL	—	MPS
Framework	TensorRT	—	OpenVINO	—	CoreML

注意：”—“表示该品牌在该层没有自己的组件。

SYCL 特殊之处： 它没有自己的 runtime，而是通过 backend plugin 使用 OpenCL RT、Level Zero 或 CUDA RT。这是唯一一个”语言层和 runtime 层完全解耦”的方案。

总结

AI/GPU 软件栈可以清晰地分为 7 层，每层解决一个明确的问题：

Hardware ISA — GPU 能执行的机器指令（厂商私有）
Driver — 硬件管理 + IR→ISA 编译器后端
Runtime — 编程接口：buffer、queue、kernel dispatch（类比 libc/JRE）
Language + IR — 写 kernel 的语言 + 编译后的平台无关字节码（类比 Java bytecode）
Operator Library — 预优化 kernel 集合 + runtime 胶水（类比 MKL/BLAS）
Graph Optimizer — 图级编译器，算子融合和内存优化
Inference Framework — 加载模型、调度算子

CUDA/OpenCL/SYCL/ROCm/oneAPI 之所以让人困惑，是因为它们每一个都横跨了多层。记住这个 7 层模型，遇到新技术名词时，先问”它在哪一层？“就能迅速定位。