一文理清 汇编、图形学API、CUDA，在完整的链路中各自的位置

引言本文要回答一个问题汇编、图形学API、CUDA它们各自在从代码到硬件的完整链路中到底处于什么位置这个问题看似简单但一旦认真追下去会发现我们对底层这个词的理解往往停在了某一层就不再继续了。我们以为脚下就是地面但脚下可能还有地下室。本文沿两条主线——CPU与GPU——逐层剥开看看每一层到底站在哪里。GPU部分本文以NVIDIA平台为具体讨论对象。NVIDIA拥有最完整的从CUDA到PTX到SASS的工具链也是目前GPU计算生态中链路最清晰、资料最丰富的平台便于完整展开讨论。其他平台AMD、Apple、Intel在架构和细节上有差异但核心结论——你和GPU之间始终隔着驱动这个黑盒——是共通的。需要提前说明本文的目的不是贬低任何一层抽象的价值。每一层包装都有它存在的工程理由。本文的目的只有一个——搞清楚每一层到底站在哪里。第一部分CPU 主线——从包装到裸金属一、裸储存内存的真相一切数据结构——std::vector、红黑树、哈希表——拆到最底下是什么是字节。地址: 0x0000 0x0001 0x0002 0x0003 ... 内容: [0x4F] [0xA3] [0x00] [0x12] ...内存不知道自己存的是整数还是字符串。它只是一个巨大的字节数组。类型是人类赋予的含义不是内存自身的属性。unsignedcharmem[8]{0x48,0x65,0x6C,0x6C,0x6F,0x00,0x00,0x00};char*as_string(char*)mem;// 看到 Helloint*as_int(int*)mem;// 看到 1819043144float*as_float(float*)mem;// 看到一个浮点数// 同一块内存不同的眼睛看到不同的东西// 内存本身从未改变变的是你解读它的方式所谓结构体不过是偏移量约定——前4字节叫x后4字节叫y字段名只是编译器帮你记住的偏移量。所谓数组不过是起始地址步长。所谓链表不过是一块内存里藏着另一块内存的地址。所有数据结构拆到底都是字节和地址的游戏。不过需要注意一点这不是世界的本质。计算机本身是人造物——人类选择用电压高低表示0和1选择用8个bit编为1个byte选择用线性地址组织内存。这些都是设计决策不是自然规律。我们往下追追到的不是自然真理而是另一群工程师的设计意图。但在这台人造机器的范围内字节就是最底层的存储单位。二、裸指针一个被过度包装的整数剥掉shared_ptr、unique_ptr、引用、句柄指针的真相极其朴素指针就是一个整数它的值是一个内存地址。解引用就是去那个地址读写数据。intx42;int*px;// p 里存了一个数字这个数字是 x 的地址*p100;// 去那个地址把那里的内容改成 100在64位机器上指针就是一个64位无符号整数指向内存这个巨大字节数组中的某个位置仅此而已。#includestdio.h#includestdint.hintmain(){intx42;int*px;uintptr_traw(uintptr_t)p;// 把指针转成整数printf(指针的本体: %lu\n,raw);// 输出一个普通的数字// 它就是一个数字// 这个数字的意思是去内存的第 raw 号格子看看return0;}所有围绕指针的智能指针、所有权语义、生命周期管理、借用检查都是在这个朴素事实之上的人为约束——为了防止人犯错而加的规矩不是因为指针本身有多复杂。三、类型系统一场善意的遮蔽CPU不认识类型。内存不认识类型。从晶体管到寄存器到总线流动的只有0和1。类型是人类在编译器层面施加的约束有类型流派C/Java/Rust... 先定义是什么再允许操作 int x 42; // 必须先声明类型 哲学先分类再使用 无类型流派汇编/Forth... 一切都是数据解读权在使用者 哲学数据就是数据类型是后加的C语言的void *是对裸内存真相的一种致敬void*pmalloc(100);// 100个字节没有类型int*a(int*)p;// 你说它是整数char*b(char*)p;// 你说它是字符double*c(double*)p;// 你说它是浮点// 内存没变变的只是你的解读类型系统是一层善意的遮蔽——它挡住了你直接面对裸字节的自由换来不容易犯错的安全。对于大规模工程来说这是值得的交换。但你需要知道它遮住了什么。四、从高级语言到机器码逐层下降沿着CPU这条线一路往下每一层到底做了什么第五层高级语言 std::vectorint v; 大量抽象RAII、模板、迭代器、异常处理 你在和概念打交道 ↓ 第四层C 语言 int *p malloc(40); 抽象变少手动内存管理但还有函数、类型、变量名 你在和操作系统接口打交道 ↓ 第三层汇编 mov eax, 42 几乎只剩命名操作码有名字寄存器有名字 你在和CPU指令打交道 ↓ 第二层机器码 B8 2A 00 00 00 纯字节CPU真正读取和执行的东西 ↓ 第一层电信号 高电压 / 低电压 不再是编程而是物理这里有一个关键分界线。从汇编到机器码是一一对应的mov eax, 42 ←→ B8 2A 00 00 00 add eax, 1 ←→ 83 C0 01 jmp loop ←→ EB F6 ret ←→ C3mov是操作码B8的名字eax是寄存器编号000的名字loop是某个具体跳转地址的名字。汇编器做的事情是查表替换——不是翻译不是优化不是重组。而从C语言到汇编呢// 你写的 C 代码intsum0;for(inti0;i100;i){sumarr[i];}// 编译器可能生成的汇编开启优化后// 向量化、循环展开、指令重排...// 你已经认不出自己写的代码了从C到汇编编译器在替你做大量决定。从汇编到机器码没有任何人替你做决定。这是本质区别。五、汇编的定位人与CPU之间的完美平衡点现在可以准确地定位汇编了往上一步C语言 出现了类型、函数、变量名、控制结构 编译器开始替你做优化决定 你写的代码 ≠ 机器执行的代码 开始失去对硬件的完全掌控 汇编这里 你写什么CPU就执行什么 1:1 对应零黑盒 完全的自由完全的掌控 同时仍然是人类可读的符号 往下一步机器码 自由度与汇编完全相同 但失去了可读性 B8 2A 00 00 00 和 mov eax, 42 是同一条指令 一个人能读一个人读着极其痛苦 再往下电信号 不再是编程而是电路汇编是最后一层人类能舒适读写的东西。它没有增加任何约束没有减少任何自由只是给机器码起了人类能读的名字。汇编是人类与CPU之间最薄的一层翻译——薄到再削一刀就不是语言了。这就是汇编在CPU链路中的位置完美的平衡点。它给你完全的自由和掌控同时让你还能作为人类去阅读和思考。CPU世界里没有比这更准确的位置了。第二部分GPU 主线——两条路同一堵墙GPU驱动之上分出两条路径应用层 / \ / \ CUDA 图形学API 计算 Vulkan/DX12/OpenGL/Metal \ / \ / NVIDIA 驱动程序黑盒 ↓ NVIDIA GPU 硬件一条面向通用计算CUDA一条面向图形渲染Vulkan/DX12/OpenGL/Metal。它们都声称自己很底层。下面分别展开。一、图形学API路线Vulkan / DX12 / OpenGL / Metal1.1 先看它上面有多少层在讨论图形学API底不底层之前必须先看清它上面站着什么应用层终端用户软件 Photoshop / After Effects / Premiere 用户点按钮不知道底下发生了什么 应用层DCC工具数字内容创作 Maya / 3DS Max / Blender / Houdini 艺术家操作场景、模型、材质 应用层游戏引擎 / 渲染引擎 Unreal Engine / Unity / Godot 封装了整个渲染管线、物理、资源管理 第三层图形学API ← 我们讨论的这一层 Vulkan / DX12 / OpenGL / Metal 第二层GPU驱动程序 NVIDIA 驱动私有实现 第一层GPU硬件 NVIDIA GPU执行 SASS 机器指令1.2 向上看图形学API是绝对的底层站在引擎和DCC工具的角度向下看图形学API就是地基一个 Unity 开发者的世界 gameObject.GetComponentRenderer().material.color Color.red; 这一行代码背后 Unity 更新材质属性 Unity 决定渲染队列 Unity 生成绘制命令 Unity 调用图形学API提交命令 NVIDIA 驱动把命令翻译成 SASS 指令 GPU 执行 从 material.color Color.red 到 GPU 执行 中间隔了整个引擎的渲染管线一个 Unreal 开发者的世界 蓝图里连几根线拖几个节点 底下是 UE 的 RHIRender Hardware Interface RHI 底下才是 Vulkan/DX12 Vulkan/DX12 底下是 NVIDIA 驱动 驱动底下才是 GPU 开发者距离图形学API隔了 蓝图 → C → RHI 三层一个 Blender / Maya 用户的世界 鼠标拖一个模型点一下渲染 底下是 Cycles / Arnold 渲染器 渲染器底下是 OptiX / CUDA 或 OpenGL / Vulkan 再底下是 NVIDIA 驱动 再底下是 GPU 用户距离图形学API隔了整个DCC软件的架构从这个角度看说图形学API很底层完全成立。它是引擎和DCC工具脚下的地基。能直接写Vulkan或DX12意味着你在直接控制渲染管线的每个阶段——自己管理命令缓冲区、自己做同步、自己分配显存、自己编译着色器。相对于在Unity里拖组件或在Maya里点按钮这已经是非常不同层次的工作了。这一点必须承认也必须尊重。1.3 向下看但它不是GPU的汇编承认了图形学API向上看的底层地位之后继续往下追。先看这三个字母Application Programming Interface Application → 应用层的 Programming → 编程用的 Interface → 接口接口。不是硬件本身是硬件前面的一道门。这个名字从第一天起就在告诉你它的定位。以NVIDIA平台为例你写的着色器代码从提交到执行经历了什么你写的 GLSL / HLSL着色器源码 ↓ 编译成 SPIR-V / DXIL中间表示标准化的字节码 ↓ NVIDIA 驱动程序再编译成 SASSNVIDIA GPU 真正的机器指令 ↓ NVIDIA GPU 执行你和GPU之间隔了两层编译加一个驱动。你从未直接向GPU发出过一条指令。显存也不在你手里// 你以为你在操作显存vkAllocateMemory(device,allocInfo,nullptr,memory);// 实际上// 你在请求NVIDIA 驱动帮你分配// 驱动决定分配在显存的哪个位置// 驱动决定什么时候搬运数据// 可能在显存也可能在系统内存// 你拿到的是一个句柄——间接的间接的引用// 你从未拿到过一个真正的显存物理地址图形学API的工作模式本质上是填表提交等待驱动执行// Vulkan 的典型工作流// 第一步填表——填一大堆描述符告诉驱动你想要什么VkGraphicsPipelineCreateInfo pipelineInfo{};pipelineInfo.sTypeVK_STRUCTURE_TYPE_GRAPHICS_PIPELINE_CREATE_INFO;pipelineInfo.stageCount2;pipelineInfo.pStagesshaderStages;pipelineInfo.pVertexInputStatevertexInputInfo;pipelineInfo.pInputAssemblyStateinputAssembly;pipelineInfo.pViewportStateviewportState;pipelineInfo.pRasterizationStaterasterizer;// ... 还有很多字段要填// 第二步提交——把填好的表格交给 NVIDIA 驱动vkCreateGraphicsPipelines(device,cache,1,pipelineInfo,nullptr,pipeline);// 第三步等——驱动拿着这些描述去办事// 你无法控制驱动具体怎么执行// 你不知道驱动做了哪些优化和重排// 你只能等结果这不是在操作GPU这是在请求GPU。你填好表格交给NVIDIA驱动这个中间人中间人替你去办事。你对最终的执行没有直接控制权。1.4 图形学API的准确定位综合向上和向下两个方向的观察图形学API在GPU世界里的位置 大致相当于 C 语言在CPU世界里的位置 CPU 高级语言 → C语言 → 汇编 → 机器码 GPU 引擎/DCC → 图形学API → 驱动 → GPU指令 向上看——它是底层毫无疑问 向下看——它下面还有驱动和GPU指令集 它是中间层不是终点图形学API的价值是巨大的——它让你能直接控制渲染管线它是引擎和DCC工具的地基它相对于上层应用是绝对的底层。但它和GPU硬件之间隔着一个你无法穿透的驱动。二、CUDA路线通用GPU计算2.1 CUDA 给了你什么以NVIDIA平台为例CUDA提供的能力相当深入直接思考线程 / warp / block 手动管理 shared memory 控制内存合并访问coalesced access 感知 SM 数量、warp 大小、寄存器数量 写 PTX 内联汇编 用 cuobjdump 查看编译出的 SASS你能感知NVIDIA GPU的架构细节并据此优化代码__global__ void kernel(float *data) { __shared__ float smem[256]; // 直接操作 shared memory int tid threadIdx.x; int gid blockIdx.x * blockDim.x threadIdx.x; smem[tid] data[gid]; // 全局显存 → 共享内存 __syncthreads(); // 手动同步 // 你在思考线程、线程块、共享内存、同步 // 这已经是 GPU 架构级别的思维 }CUDA甚至允许你写PTX——NVIDIA GPU的中间表示语言__device__ int my_add(int a, int b) { int result; asm(add.s32 %0, %1, %2; : r(result) : r(a), r(b)); return result; }从计算的角度CUDA已经是NVIDIA GPU编程能力的天花板了。2.2 CUDA 向上看比图形学API更深一步CUDA vs 图形学API 图形学API 你在填描述符、提交命令缓冲区 你在描述我想渲染什么 你不直接思考 GPU 的硬件结构 CUDA 你在直接编排线程、管理内存层级 你在思考GPU 硬件怎么运转 你能感知 warp、SM、shared memory、寄存器这意味着CUDA在GPU链路上比图形学API更深一层引擎/DCC → 图形学API → CUDA → 驱动 → GPU指令 ↑ 你在这里能感知更多的硬件细节2.3 CUDA 向下看链路仍然不透明但是CUDA的链路仍然经过两道处理你写的 CUDA C/C ↓ nvcc 编译成 PTXNVIDIA 的中间表示还不是最终指令 ↓ NVIDIA 驱动再编译成 SASSGPU 真正的机器指令 ↓ NVIDIA GPU 执行即使你直接写了PTXNVIDIA驱动仍然会把PTX再编译一次生成最终的SASS。驱动可能重排你的指令可能改变寄存器分配可能做你不知道的优化。你写的代码 ≠ GPU最终执行的代码。这和CPU汇编形成了鲜明的对比CPU 汇编 mov eax, 42 → B8 2A 00 00 00 → CPU 执行 你写的 机器执行的 1:1 对应零黑盒 汇编器只做查表替换不做任何优化 CUDA即使写到PTX级别 你的PTX → NVIDIA 驱动再编译 → SASS → GPU 执行 你写的 ≠ 机器执行的 中间还有一个黑盒在帮你做决定2.4 CUDA的准确定位CUDA 从能力上说 极其强大 几乎是 NVIDIA GPU 计算的天花板 比图形学API更深一层 CUDA 从链路透明度上说 不具备CPU汇编的地位 CPU汇编的地位在于1:1——你写什么CPU就执行什么 CUDA做不到这一点——你和SASS之间还隔着驱动编译器三、驱动两条路汇合处的那堵墙无论走图形学API路线还是CUDA路线都会在同一个地方被拦住——NVIDIA驱动程序。图形学API CUDA \ / \ / \ / NVIDIA 驱动程序黑盒 ↓ NVIDIA GPU 硬件驱动到底做了什么指令编译把 SPIR-V / PTX 编译成 SASSGPU 真正的机器码 内存管理决定数据放在显存的哪个位置何时搬运 指令调度决定多个任务如何在 GPU 上并行执行 功耗管理根据负载调整 GPU 频率和电压 错误恢复GPU 挂了尝试重置 硬件适配同一个API调用在不同型号的 NVIDIA GPU 上走不同路径 性能优化指令重排、寄存器分配、占用率优化它是黑盒的原因1. 商业机密——SASS 指令集是 NVIDIA 的核心竞争力 2. 硬件差异——同一代驱动要支持几十种不同型号的 GPU 3. 迭代速度——GPU 架构每一两年就换一代Turing→Ampere→Ada Lovelace→Blackwell 4. 优化策略——驱动的编译优化是 NVIDIA 的核心技术壁垒这堵墙意味着你永远无法确切知道你的代码在 NVIDIA GPU 上最终变成了什么指令。你可以用NVIDIA Nsight或cuobjdump去观察编译出的SASS但你无法控制编译过程本身。四、GPU为什么不给你汇编级的自由对比CPU的世界CPU 的世界 指令集完全公开 Intel 发布 x86 手册几千页每条指令都有文档 ARM 发布 Architecture Reference Manual RISC-V 甚至是开源指令集 汇编器做一一对应的翻译 你随时可以直接写机器码 你和CPU之间可以做到零距离 NVIDIA GPU 的世界 SASS 指令集没有官方完整文档 每一代架构的 SASS 都在变 Turing 的 SASS ≠ Ampere 的 SASS ≠ Ada 的 SASS 驱动是黑盒 NVIDIA 有意在你和 GPU 之间放置中间层NVIDIA GPU真正的汇编——SASS——长这样/*0000*/ MOV R1, c[0x0][0x20] /*0010*/ S2R R0, SR_CTAID.X /*0020*/ S2R R2, SR_TID.X /*0030*/ IMAD R0, R0, c[0x0][0x28], R2 /*0040*/ ISETP.GE.AND P0, PT, R0, c[0x0][0x158], PT /*0050*/ P0 EXIT无论你写Vulkan还是CUDA你从来没有直接写过这些指令。它们是NVIDIA驱动编译器的输出不是你的输入。你可以观察它们但你不能直接编写它们并交给GPU执行——NVIDIA没有给你这个入口。在CPU的世界里你可以从高级语言一路走到机器码每一层都可以亲手触碰。在NVIDIA GPU的世界里你最多走到CUDA/PTX这一层然后就被驱动挡住了。第三部分全景对照一、两条链路的完整展开══════════ CPU 链路 ══════════ ═══════════ GPU 链路NVIDIA═══════════ 应用软件 终端用户软件 Office/浏览器/游戏 Photoshop/Premiere ↓ ↓ 框架/运行时 DCC工具 / 引擎 Qt/.NET/Electron Maya/Blender/UE/Unity ↓ ↓ 高级语言 ┌──────┴──────┐ C/Rust/Java ↓ ↓ ↓ 图形学API CUDA C 语言 Vulkan/DX12 CUDA C ↓ OpenGL/Metal ↓ ↓ PTX汇编 ──── 分界线 ──── 着色器语言 中间表示 上面编译器替你决定 GLSL/HLSL/MSL ↓ 下面你说了算 ↓ ↓ ────────────── ┌────┴─────────────┘ ↓ ↓ 汇编语言 NVIDIA驱动黑盒 mov eax, 42 指令编译/内存管理 ↓ ← 1:1 零黑盒 调度优化/硬件适配 机器码 ↓ B8 2A 00 00 00 SASSGPU机器码 ↓ ↓ CPU 执行 NVIDIA GPU 执行二、三者的定位总结┌──────────────────────────────────────────────────────────────────────┐ │ │ │ 汇编CPU链路 │ │ ───────────── │ │ 位置人类与CPU之间最薄的一层翻译 │ │ 核心特征1:1 对应机器码零黑盒 │ │ 你写什么CPU就执行什么 │ │ 往上一步就有编译器替你做决定往下一步就不是人类可读的语言了 │ │ → 完美的平衡点 │ │ │ │ 图形学APIGPU链路·渲染路线 │ │ ────────────────────────── │ │ 位置NVIDIA驱动之上引擎/DCC工具之下 │ │ 向上看是引擎和DCC工具的地基绝对的底层 │ │ 向下看只是一个接口下面还有驱动和GPU指令集 │ │ 工作模式填表 → 提交 → 等驱动执行 │ │ → 相对底层但不是GPU的汇编 │ │ │ │ CUDAGPU链路·计算路线 │ │ ──────────────────── │ │ 位置比图形学API更深一层但仍在NVIDIA驱动之上 │ │ 能力感知GPU架构可写PTX几乎是NVIDIA GPU计算的天花板 │ │ 但PTX仍经驱动再编译为SASS你写的≠GPU执行的 │ │ → 能力极强但链路不具备CPU汇编的1:1透明度 │ │ │ │ 三者共同点 │ │ 都有巨大的价值都是各自领域内不可替代的工具 │ │ 区别只在于在从代码到硬件的完整链路上它们各自站在不同的位置 │ │ │ └──────────────────────────────────────────────────────────────────────┘结语汇编站在CPU链路的最底层可编程位置——它和机器码一一对应中间零黑盒你写什么CPU就执行什么。它是人类与CPU之间最薄的翻译层。图形学API以NVIDIA平台上的Vulkan/DX12/OpenGL为例站在GPU链路的中间——向上看它是引擎和DCC工具的地基绝对的底层向下看它下面还有NVIDIA驱动和SASS指令集它只是一个接口。CUDA站在GPU链路中比图形学API更深的位置——它让你直接思考NVIDIA GPU的架构甚至可以写PTX。但PTX仍然经过NVIDIA驱动再编译才变成SASS链路不透明。GPU世界之所以没有等价于CPU汇编的东西是因为NVIDIA以及其他GPU厂商有意不给你那个入口。指令集是商业机密驱动是黑盒每代架构都在变。这是商业决策不是技术限制。认清每一层的位置不是为了否定任何一层的价值——每一层包装都有它存在的工程理由。但你至少应该知道你站在哪里你脚下还有什么以及那些替你做决定的黑盒到底遮住了什么。