企业官网建设流程全解析

OpenDLSS-NR是 NVIDIA DLSS 5 神经渲染（Neural Rendering）网络的 Vulkan 重实现，与原实现逐位一致（bit-exact）：同样的 71 块 Swin / ViT 网络、FP8（E4M3）激活、141 MiB 的权重。它最大的门槛不在编译，而在模型目录——权重不是常见的 ONNX / safetensors，而是一堆按 tensor core 的 MMA 碎片顺序打包的 E4M3 字节。本文带你读懂 manifest.json 的结构、153 条 tensor 记录的组织方式，以及 packedWeightIndex 背后的 E4M3 MMA 碎片打包布局。

同一帧画面，左关右开神经渲染（WebGPU 移植版，2048x1152）。

模型目录长什么样：manifest.json 三要素

仓库不提供权重，你需要自备一个模型目录（默认在 README 旁的models\nr，也可用--model <dir>指定）。目录的核心是manifest.json，由 src/nr_model.cpp 的Model构造函数解析：

一句话概括：stage 是磁盘上的物理文件，tensor 是逻辑切片。加载时 stage 被整块读入内存（L78-L89），153 条 tensor 记录则各自指向对应 stage 的字节区间，零拷贝完成映射。

💡 整个模型的加载耗时约几秒，其中大头是 SHA-256 校验——这恰好说明 141 MiB 权重的完整性是加载流程的一部分，而非可选步骤。

一条记录里有什么：FusedLayout的权重拼接顺序

一条 tensor 记录不是"一个矩阵"，而是一整层的所有权重首尾拼接：FFN 权重、逐通道 f16 缩放向量、注意力先验、逐头 f32 缩放……顺序由 FusedLayout 定义。以 C 通道的窗口块为例（摘自 docs/weights.md）：

官方文档给出了一条漂亮的自检逻辑：FusedLayout里每个偏移都是前面形状之和，这个和恰好精确复现每条 tensor 的byteLength——布局对不对，算一遍就知道。

不同块还有各自的变体：block 0 在 W2 后插入16→32输入适配器（preFusedLayout）；每个解码器阶段首块插入上采样矩阵（upsampleFusedLayout）；block 70 追加32→4的 f16 头（postFusedLayout）；512 通道块拆成 4 条 layer 记录；ViT 块则把每头缩放放在 qkv 权重之前。

核心：E4M3 权重为什么按"碎片"打包

这是全文的重点。docs/weights.md 开篇就点破：权重不是按矩阵存的，而是按 MMA 碎片存的。

网络里每个 FP8 GEMM 都是一串16x16x32的 E4M3×E4M3→f16 协同矩阵乘法，对应硬件指令mma.sync.aligned.m16n8k32.row.col.f16.e4m3.e4m3.f16。于是最"顺路"的存储顺序就是 B 操作数在寄存器里的碎片顺序：

• 最内层 512 字节 = 一对m16n8k32B 碎片：16 列 × 32 个 k，分布在 32 条 lane 上，每条 lane 16 字节（两个 8 列半区 × 两个 16-k 半区 × 4 字节）；
• 往上是nGroup（512 B）、nHalf（2048 B）、nTile（4096 B）、kTile（N·32B）逐级卷叠。

字节级索引公式在 packedWeightIndex 中只有几行位运算：

lane = (n % 8)·4 + (k % 16)/4 byteInLane = ((n % 16)/8)·8 + ((k % 32)/16)·4 + (k % 4)

直观理解：lane 编号决定"哪条 GPU 线程"，lane 内 16 字节决定"这个线程拿到的 4 个 k × 4 个 n"——和 tensor core 的取数路径完全对齐，零重排即可喂给 MMA。

加载时，主机端会遍历这个碎片布局，反向摊平成普通矩阵再上传 GPU（fp8MatrixBytes）：

• 默认[K/batchK][batchK/32][N][32]（k32-tile-major），保证一个 B tile 是连续的 16 字节对齐块，配合协同矩阵加载；
• tileMajor = false时改为[K/batchK][N][batchK]（N-major），供单次乘法的 B 操作数使用。

一个容易忽略的置换：chained activation index

每个 FP8 GEMM 的 A 操作数（激活）在每个 32 通道组内还有一个位旋转置换（索引位 1..3 轮转：0 1 2 3 4 5 6 7 8 …→0 1 8 9 2 3 10 11 4 5 12 …）。

运行时对激活做置换太贵，于是 OpenDLSS-NR 把逆置换直接烘进了权重的 K 行：加载时对每个权重矩阵的 k 行做 inversePackedInputIndex 重排，内核就能按自然通道序加载 A。这在数学上成立的前提是：置换始终不越出每个 16 个乘积的小组——正是 f16 累加的分组边界。构造函数里还有一条断言式检查：置换必须可逆，否则直接抛错。

另外两种布局：f16 矩阵与注意力先验

缓存与生命周期：重排只做一次

所有重排在加载期完成后就固化为 device-local 缓冲，按(tensor, offset, shape, batchK, layout)在Model::matrices_中缓存。重录整张图时重新请求的是同样的 key，拿到的是同样的缓冲；Model析构时统一销毁，没有悬垂资源。这意味着141 MiB 的碎片解包与 SHA 校验只发生在加载那几秒，帧循环里拿到的都是现成的连续矩阵。

下一步：想继续深入读什么

跑通后可以用dlss5vk.exe bench --model <dir> --width 768 --height 768看性能，用parity --model <dir> --fixture <dir>对原始实现的录制捕获做逐位比对——权重读对了，一切才谈得上"bit-exact"。