2019 年我在 GitHub 上 fork 过一个仙剑 DOS 版的仓库，就是个 DOSBox 打包，下载下来跑模拟器。这个方案我一直不太满意，游戏数据只有几十 MB，外面却要再套一层 DOS 环境，画面是 320×200 直接拉伸。
这两天我把仓库彻底重做了一遍：引擎从 SDLPAL 的 C 代码完整移植到 Rust，两万四千行，直接在 macOS、Linux、Windows 上运行；同一份代码编译成 wasm 跑在浏览器里，在线试玩；输出 1280×720，原版 320×200 的游戏帧由一个 FP16 神经网络实时放大。开发花了三天，7 月 17 日移植，18 日超分，19 日调优。这篇按三块记录：移植、超分、kernel 调优。
Rust 移植
移植的参考是 SDLPAL，PAL_CLASSIC 经典模式，只保留 DOS 版代码路径。顺序自底向上，先数据解压和字库，再引擎和音频，最后脚本解释器、UI 和战斗。仙剑的游戏逻辑几乎全在脚本里，剧情、对话、触发器、进战斗都是 opcode 驱动的，近三千行的 script.rs 是全仓库最大的文件，也是最花时间的部分。
另一大块工夫花在验证上：
- YJ_1 解压和 C 实现逐字节对比，全部 1159 个压缩块一致。
- OPL/RIX 音乐：20 首曲子各跑 30 秒，和 C++ 原实现逐字节一致。opl.rs 刻意照搬 C 的循环和查表，怕换写法换出 bug。
- 97 个单元测试加 4 个端到端测试：比如用脚本 opcode 触发一场战斗，再拿同一个随机种子直接打一遍，经验和金钱必须逐 bit 相同。
当然啦，这个移植是我让 Claude 自动完成的，方法在之前的博文里也多次提过，比如 mihomo 移植那篇，这里只补一句：移植这种有标准答案的工作，AI 产出的上限很高，前提是验证做在它前面。
AI 实时超分
320×200 放到今天的屏幕上，nearest neighbor 放大全是方块，xBR 和 Anime4K 能把边缘做锐利但补不出细节。浏览器版我先上了这两个滤镜，随后把目标换成真正的神经网络：Real-ESRGAN 的 animevideov3 模型，专门处理动画和游戏画面，结构是 SRVGGNetCompact，18 层 3×3 卷积，64 通道，FP16 权重打包后 1.19 MB。
推理没有走 ONNX Runtime Web，整个网络手写成一个 shader module，我管它叫 mega kernel：全部权重预装进一个 GPU buffer，一帧 18 个 dispatch 录进一次 command submit，每层的权重偏移走一个带 256 字节动态偏移的 uniform buffer。激活值用 NHWC 排布的 vec4<f16>，两个 8 MB 的 storage buffer 来回 ping-pong。最后一层 conv_last 把 4 倍 pixel shuffle、nearest 上采样残差、clamp 全部融合进去，直接写出 1280×800 的 rgba8unorm 纹理。权重用个 Python 脚本从 ONNX 模型打包，38,736 个 mat4x4，带 round-trip 自检。
浏览器里有个别扭的约束：游戏画面的 canvas 已经占了 WebGL2 context，神经网络只能另起一个 WebGPU canvas 盖在上面。帧率对不齐就丢中间帧，网络跑 40 到 50 ms 一帧的时候引擎不等它。任何一步失败，降级回 xBR。
原生版复用同一份权重，.mega.bin 用 include_bytes! 嵌进二进制，wgpu 设备直接借 pixels 的，Metal 和 Vulkan 上跑同一套 dispatch 顺序。浏览器和原生用完全相同的权重和 kernel，这是我给自己定的约束，两边可以随时对拍。
精度的参照是 onnxruntime-web 的 FP16 推理：最大差 2/255，平均 0.11，偏差超过 2 的子像素一个都没有。误差来自累加顺序，视觉上不可见。代价是需要 shader-f16 这个 WebGPU feature，不支持的设备上浏览器版降级到 xBR，原生版退回 nearest。
WebGPU kernel 优化
初版 mega kernel 在 M4 上跑 48 ms 中位（纯 f16 累加；f32 累加要 89 ms），折合 21 到 25 fps，能用，但离 60 fps 很远。16 层 conv_mid 占了大约九成时间，先拿它开刀。
调优继续交给 AI 自动完成。可惜 Fable5 和 GPT5.6sol 都拒绝给我做 DL kernel 优化，只能上 Kimi K3 了。先写了一个 sweep 页面（nn-tune.html），21 个带名字的变体，GPU timestamp 计时，每个变体先把全网络输出读回内存和基线逐字节对比，一致才认它的耗时。sweep 页面的实测结果（M4，全网络中位耗时）：
| 变体 | 耗时 |
|---|---|
| 基线 8x8x1，64 线程，每线程 16 个累加器 | 54.4 ms |
| channel split=2，128 线程，8 累加器 | 30.2 ms |
| split=4 加每线程 2 像素，256 线程 | 26.5 ms（采用） |
| 权重放进 workgroup memory | 42.5 ms |
| 每线程 2 像素但不做 split | 61.2 ms |
| 手动展开 in-channel 循环 | 无收益，编译器已经展开过了 |
结论有两个。一是 occupancy 比共享内存重要：channel split 把每线程累加器从 16 压到 4，驻留线程数翻四倍，收益立竿见影；权重放 workgroup memory 反而输了，加载和同步得不偿失。二是每线程多算像素必须和 split 配合，单独上 2 像素会 register spill，比基线还慢。
最终配置 @workgroup_size(8, 8, 4)，18×10 的 halo tile，23 KB workgroup 存储。这个 23 KB 得多说一句：WebGPU 只保证 16 KB，23 KB 要请求更高的 limit，拿不到的机器退到 split=2 的 12.8 KB 配置，仍有 1.8 倍加速。
原生版把 kernel 挪到 wgpu，从 187 ms 一路修到 27 ms，主要是三件事：
- naga 默认给 workgroup memory 做零初始化，23 KB 的 tile 只用一个线程填，这一项就花掉约 110 ms。设 zero_initialize_workgroup_memory: false 解决。
- naga 的强制 loop bounding 挡住编译器展开循环，卷积链慢 2.5 倍。换 create_shader_module_trusted，关掉 bounds check 和 loop bounding。
- conv_last 同样改成 channel split 的 8x8x4，6.0 ms 降到 1.6 ms，输出逐字节一致。
之后把 conv_last 的 split 移植回浏览器版，Dawn 上耗时持平（26.48 对 26.54 ms），但浏览器和原生从此跑同一份 kernel。我也试过手写 MSL 绕过 naga，没有收益，WGSL 留下。顺带说下计时方法：Metal 的 pass 边界时间戳在这台机器上不可靠，读数能小到亚毫秒，最后全部改 wall clock，15 轮每轮 4 帧取中位和最小值，每个变体同样先过逐字节对比再谈耗时。
结语
这次重制几乎没动美术。原版 320×200 的手绘像素画，线条干净，色块大，正好是这类模型最拿手的输入。一个像素没重画，超分补出来的细节我个人觉得明显好过 xBR。全平台移植加超分加调优三天做完，真正费劲的是移植的逐字节验证和 kernel 的调参，超分本身反倒是现成模型套上去就能用。
也想借这个例子给游戏厂商提个建议：不少当年的好作品卡在过时的分辨率和平台上，新一代玩家已经玩不到了，按这套思路重制一遍成本并不高，能让现在的小孩也玩上当年的优秀作品。