翻译器的由来

系列gdev-masterNVIDIA/nouveau 用户态 GPGPU 运行时从 C/C 到 Rust 的移植工程上一篇《免疫的来由》收尾了慢带那条线。这一篇回到移植工程的主线产线——程序怎么变成 GPU 上能跑的东西。当时它还有两处要靠别人机器指令文本要人手写汇编到中间格式的翻译靠官方工具链。今天产线长这样每一段都是仓库里自己的代码与 2012 年的 C 原版逐字节等价。本篇讲这段路的三件事手工造的程序在真显卡上读回了 42真编译器的程序让 C 原版和 Rust 移植版一起卡死在同一处以及最后一段依赖——PTX 翻译器——被替换掉的全过程包括评估为什么先别做。先给结论问题结论依据手写的 GPU 机器指令能不能在 GTX 580 上跑通并读回数值能。程序往结果里写 42目标机读回got42 expected42 PASS两次运行都过set42两跑真编译器产出的、带常量的程序能不能跑不能。C 原版和 Rust 移植版卡死在同一处——重定位的应用代码 2012 年就是空循环两侧源码对照PTX→SASS 翻译器能不能 Rust 化并逐字节一致能。评估先别做五期后收口cargo 测试 116 个双侧全绿冻结判据 4 套 56 区段失败方式一致吗一致。8 个 2011 年的测试件两侧在同一层以同类断言失败定位可对到行号负例全套回放是不是全量指令集都验证了不是。逐字节的范围 冻结语料覆盖的支持集支持集外与 C 同点同方式 fail-fast判据定义在案一、手工造一个 GPU 程序读回 42这套工程的验收标准是行为与原版逐字节一致加载 GPU 程序的部分原本靠现成程序文件测试——但现成文件都是别人编译好的整条链在真显卡上走没走通过始终缺一锤。编译器是个麻烦GTX 580 是 2010 年的卡官方编译器对它的支持停在 2013 年前后现代电脑装不上。好在仓库里留着原作者的小工具sass2cubin吃手写的汇编文本SASS这一代显卡的机器指令打包成cubin。理论上可以不依赖编译器自己把指令一条条写出来。流水线前半截在 2012 年系统的目标机上官方工具只在那里有probe.cu一行本体out[threadIdx.x] 42 → nvcc -archsm_20 -cubin # 官方编译器只借它产出标准格式 → nvdisasm -c # 官方反汇编器看真指令长什么样 → 手工整理成 set42.sass # 照着改写成汇编器能吃的文本 → sass2cubin -sm_20 -64 # 2012 年的小工具打包成 cubin产物 1,202 字节。三个坑都是现场踩出来的反汇编器挑食直接反汇编仓里现成的 cubin 报不含架构信息换标准格式就正常寄存器名差一个点特殊寄存器认SR_Tid_X写成SR_TID.X直接拒绝64 位参数要拆两半传两次 32 位接口低半和高半落在参数区固定的两个位置。最难的从来不是指令本身——算出线程号、取参数指针、写 42、退出就这几条——而是让每个工具在同一套约定里说话架构代号、指针宽度、参数布局一个对不上下游就静默出错。然后加载、启动、读结果用的是移植后的 Rust 库got42 expected42 PASS两跑皆过。二、真编译器的程序连失败都一模一样第二组实验换了真编译器产出的、稍微真实的程序一个常量、一个全局变量、一个把两者相加写出的函数。这种程序里编译器对全局变量的地址先留空位指望加载器在加载时把真实地址填上——这就是重定位先占位、后填数的那类修补。结果加载每一步都报成功启动也返回成功然后卡死在等结果那一步直到超时。排查把范围收得极窄C 原版一模一样地卡死而且原因在源码里看得见。解析到重定位表时原作者的记档是gdev_cuda.c:521caseSHT_REL:/* relocatable: not sure if nvcc uses it... */而应用它的地方是个循环——循环体里只有注释gdev_cuda.c:634-645for(sym_entry(symbol_entry_t*)nvrel;(void*)sym_entry(void*)nvrelsheads[nvrel_idx].sh_size;sym_entry){/* ……原本该做的事全部以注释形态存在 */}打算做的一切都以注释写好了只是没写成代码。Rust 侧逐行对译C 是什么样它就是什么样——同样空的迭代注释里还专门记了忠实保留。这是一个 2012 年就存在的能力缺口不是移植引出来的 bug。而这是移植工程想要的结果之一行为等价的意思是成功要等价、失败也要等价——原版以什么方式失败移植版就以什么方式失败一步不差。假如移植版好心把重定位实现了两侧反而出现分歧验收标准就被破坏了。裁决随之而定补齐重定位不立项真有需求先定 C 原版怎么改不能只改 Rust 一侧。三、最后一段依赖到这里产线还剩最后一段要靠官方工具程序文本到汇编文本的翻译。现代工具链对这张 2010 年的卡早已停止支持但仓库里有个 2011 年的翻译器ptx2sass吃 PTX 文本NVIDIA 的中间汇编现代编译器仍以它为输出吐 SASS 文本机器指令的可读形式。立项前照例评估两个坏消息它不是一个工具。编译时拖出的是一整棵 175k 行的 Ocelot 源码树12 个模块全编进一个静态库工具只是树梢一根枝。它是另一棵树。此前移植过的 Ocelot 运行时部分与它同源不同版本diff -rq实测代码零复用此前对 Ocelot 的了解帮不上忙。四、把翻译 175k 行变成回放 56 段字节真正动手前规模缩了三次水每次都有实测背书175k 行的树编译可达的只有工具级闭包——从main出发真正会被执行的代码约 25-40k 行核心翻译器本体只有 1,490 行。闭包里还有不可观察的部分——比如一个 581 行的调试输出函数后面讲它为什么最终一行都没移植。每一层 Rust 代码的验收不靠看起来对靠和 C 原版比字节。第三条是方法的支柱。C 原版构建成功五个编译适配补丁全部零行为变化、全部脚本化它成了每一层的证人喂它输入把它内部关键状态的文本快照冻结成 golden 文件跑两遍确认完全一致才冻结Rust 侧重建同样的快照逐字节比对。冻结了四套层冻结内容区段数基础设施命令行解析/文本格式化/版本号 SASS 指令表示25PTX 前端解析器逐语句的内部状态流 4 个白名单负例14翻译核心控制流块/寄存器编号/控制树形状/翻译后指令2CLI 整体退出码/断言行/输出文件/help 全文正例加负例15五、翻译器的里子一个物理寄存器伺候三个变量对外行说一句这翻译器在干什么PTX 里变量叫虚拟寄存器数量不限GPU 上真实的物理寄存器只有几十个。翻译器的核心工作之一是分配——让有限的物理寄存器接力伺候无限多的虚拟寄存器谁的生命周期结束谁的槽位立刻让给下一个。这个 2011 年的分配器用了个朴素的办法先从头到尾扫一遍给每个虚拟寄存器记下第几行用到、第几行最后一次用到的行区间翻译时谁在区间内谁占槽区间过了就释放。看一眼冻结的 golden 输出就懂——同一个物理寄存器 R6开头 7 条指令里接力伺候了三个不同的虚拟寄存器/*0020*/ MOV32I R6, 0x2; ← 虚拟寄存器 %r7 占用 R6 /*0028*/ MOV32I R6, 0x1; ← %r7 的区间已过R6 让给 %r6 /*0030*/ IADD R6, R4, R5; ← %r6 的区间也已过R6 再让给 %r3这种看起来像 bug 的正确行为到处都是分配器的槽位释放表在极端情况下会误删易主后的新占用C 原版如此跳转目标在中间表示里被改写成内部块名函数返回被翻译成退出……移植的原则先证明它是原版行为然后保留。本轮实测归档的这类怪癖前后共二十余条。六、三个真发现1. 手推两次全错的算法冻结一次对齐。翻译器依赖一个控制树分析把控制流图归并成一棵结构树决定指令的最终顺序。它的内部集合用的是 C 无序集合——遍历顺序取决于对象在内存里的地址哈希。手工在纸上推演这个顺序推了两次两次都错区域编号、成对结构都岔了。放弃手推写探针程序把 C 的真实树形冻结下来——Rust 按插入序模拟一次对齐连树节点编号都分毫不差。教训很便宜也很值钱凡是实现定义的顺序纸笔推演都是自我安慰实测冻结才是证据。2. 分层全绿全链翻车。有个立即数格式的数据C 里靠联合体存储——同一块内存按浮点写、可按整数读。Rust 没有这种一写多读移植时把每个面单独存字段。解析层的判据先建逐字节比对解析状态流全绿。几天后全链判据从 PTX 文本一路比到 SASS 文本一开翻车——浮点位型的立即数解析层绿到输出变成 0。查下来是联合体有三个可读面移植只补了两个。这道门的价值在于每一层各自全绿加起来不等于全链绿。分层判据负责定位全链判据负责兜底缺一不可。3. 失败都要对上号。仓库里躺着 8 个 2011 年的测试文件全是原版跑不通的——覆盖了当时未实现的指令。移植验收把它们也收进来要求 Rust 版以同样的方式失败。实测比预期还细8 个文件全部在翻译层终止而非解析层终止时的断言定位文件名、行号、函数签名逐一与 C 原版对上。有两个多函数文件C 与 Rust 首个失败点不同——深挖发现是 C 用了依赖内存地址的遍历顺序原则上不可复现判据遂如实分层单函数文件比到行号多函数文件只要求同一层、同类断言。七、能说到哪不能说到哪能说一例手工 kernel 的完整链——手写汇编、打包、加载、上卡、读回数值——在真显卡上走通带重定位程序的失败模式双侧一致且已定性PTX 文本 → SASS 文本翻译器在冻结语料覆盖的支持集内算术/浮点/谓词/分支/访存/参数传递与 C 原版逐字节一致含退出码与失败方式。产线最后一段官方工具依赖被自有代码替代。不能说“全量”。一例不是矩阵2 个正例 8 个负例是当前全部证人更多指令形态的语料扩充排了触发式计划不加语料就宣称全量是撒谎。32 位 cubin 两侧同崩都按 64 位解析明确不支持。边界多函数模块的内部遍历序依赖内存地址哈希原则上不可复现判据已如实分层翻译器的三个构建适配补丁只存在于 oracle 环境仓库源码未动一字。一句话总结手工汇编上卡读回 42真编译器的程序连失败都与原版同步最后一段工具依赖被逐字节一致的 Rust 翻译器替代——让每一层都有 C 做证人字节为凭。