Rubin架构FP4 GEMM深度解析：低精度推理如何成为GPU新赛道

最近圈子里讨论最热闹的下一代 GPU毫无疑问是 NVIDIA 的 Vera Rubin 平台。我翻了不少公开资料和技术拆解发现大家对 Rubin 最关心的点基本都落在 FP4 GEMM 上。FP4 这个精度从 Blackwell 开始正式走上 Tensor Core 的舞台到了 Rubin 这里它不再是边角料而是直接被当成性能指标的“扛把子”。那么 Rubin 的 FP4 GEMM 到底有什么特殊的为什么它能成为下一代推理优化的核心赛道如果你是做模型推理优化、量化部署或者单纯对 GPU 架构感兴趣这篇“概览”型的文章应该能帮你把零散信息串起来。先说明一下本文不是官方文档很多地方是结合公开信息和架构演进的合理推断我会把这些推断明确标出来大家辩证着看。1. 为什么 FP4 GEMM 成了 Rubin 的代名词1.1 先把 FP4 和 GEMM 这两个黑话掰开FP4 就是 4 位浮点数通常采用 E2M1 格式1 位符号、2 位指数、1 位尾数。相比于 FP8E4M3 或 E5M2和 BF16FP4 的位宽只有它们的二分之一或四分之一能表示的数据范围和精度自然更有限。GEMM 是 General Matrix Multiplication也就是通用矩阵乘。在深度学习里线性层、卷积、注意力机制中的投影和打分操作本质上都能归结成矩阵乘。所谓 FP4 GEMM就是让两个 4 位浮点矩阵直接相乘而不是像以前那样先把数据转回 FP16 再算。但为什么 FP4 和 GEMM 要放在一起说因为神经网络里绝大多数计算量都集中在 GEMM。把 GEMM 的输入精度砍到 4 位等于让同一个 Tensor Core 时钟周期内能处理的元素数量翻倍甚至翻四倍。按位宽估算FP4 乘法器的面积、功耗都比 FP16 小得多因此单位芯片面积上可以塞进更多的计算单元。Rubin 把 FP4 GEMM 作为宣传亮点意味着它在硬件层面不只是“支持一下”而是专门为 4 位矩阵乘优化了微架构。这里还要澄清一个常见误区FP4 GEMM 不等于“结果也存成 FP4”。矩阵输入虽然是 4 位但内部累加器通常保留 FP32 精度的中间结果最后再通过缩放因子和舍入转换才写回低精度。也就是说FP4 GEMM 是一个“输入低精度、输出高精度累加”的过程。这个设计思路从 FP16 时代一直延续到现在只是输入位宽一步步从 16 降到 8 再降到 4。1.2 量化的大趋势AI 推理正在被低位宽重塑现代大模型动辄几百亿参数推理时最大的痛点不是“算不动”而是“显存装不下、带宽跟不上”。GPT 这类自回归模型在 decode 阶段一个 token 一个 token 地生成每次都要把整个模型权重从显存搬到计算单元。权重位宽减一半带宽占用就减一半缓存容量也能塞下更多权重。所以近几年 INT8 量化、INT4 量化早已是工程标配。FP4 相比 INT4 有一个明显的优势指数位带来了动态范围。权重和激活值经常会出现少量特别大或特别小的离群值纯整数的 INT4 遇到这些值很容易直接溢出或饱和浮点格式则能通过指数把数值范围拉得更宽。虽然 FP4 的精度仍然很粗糙但在配合缩放因子的情况下经验上很多模型都能承受这种损失尤其是加入量化感知训练后精度可能逼近 FP16 基线。这也就解释了为什么 Blackwell 开始引入 FP4 Tensor CoreRubin 又进一步把它做强。推理侧的性价比模型已经从“能用 FP16 就行”变成了“能上 FP4 就上 FP4”。当大家都把模型量化到 4 位时GEMM 的性能就直接决定了整个推理系统的吞吐上限。1.3 算力参数表的暗示FP8 是训练主菜FP4 是推理甜品NVIDIA 历代 GPU 的规格表里峰值算力通常按 FP32、FP16、FP8、FP4 等分列。Hopper H100 只提 FP8 Tensor CoreBlackwell B200 把 FP4 加进了指标中而 Rubin 的官方预告里FP4 被提到了一个非常显眼的位置。这不是简单的市场宣传而是明确告诉你Tensor Core 的物理设计已经预留了 FP4 的原生路径。我个人的理解是FP8 依然会是训练阶段的主力因为反向传播对梯度精度要求更高FP4 很难撑起大批量训练。但在推理阶段模型权重已经固定可以通过离线量化把权重压到 FP4激活值也能节省大量带宽。Rubin 的定位很可能是“推理专用算力大幅提升同时保持 FP8 训练能力”。如果你看到未来发布规格中 FP4 稀疏算力是 FP8 的两倍以上完全不用惊讶这正是架构层面的有意取舍。2. Rubin 架构怎么为 FP4 GEMM 铺路2.1 Vera Rubin 平台概览Vera Rubin 是 NVIDIA 下一代计算平台的名字包含 Vera CPU 和 Rubin GPU 两部分。Rubin GPU 采用双裸片chiplet设计两个计算 die 之间通过 NV-HBI 高带宽接口互联搭配 HBM4 显存。整个平台还引入了 NVLink Hub用于多 GPU 之间的互联拓扑。这些听起来和 FP4 GEMM 没直接关系但实际上环环相扣。GPU 从一个架构走到另一个架构最大瓶颈往往是“喂给计算单元的数据不够快”。Rubin 之所以把 HBM 从 HBM3E 升到 HBM4把互联从传统 PCB 走线改成更先进的基板封装就是为了让每个时钟周期都能搬运足够多的低精度数据。FP4 GEMM 的计算密度异常高如果显存带宽不够芯片就会大量空转浪费时间在等数据上。这里值得提一句Vera Rubin 的命名来自天文学家 Vera Rubin这也暗示 NVIDIA 想强调“探索未知”的调性。对我们做工程的人来说名字不重要重要的是它背后的物理设计是否能真正兑现超过上一代的推理性能。2.2 Tensor CoreFP4 乘法器密度怎么玩Tensor Core 是目前 GPU 上最适合矩阵乘的专用单元。从 Volta 时代起它就固定按小 tile 进行矩阵乘。一般来说一个 Tensor Core 内部由多个 4x4 或 8x8 的乘累加阵列组成每个周期算出一个子矩阵块的部分结果。Hopper 的 Tensor Core 针对 FP8 做了专门设计Blackwell 则扩展到 FP4。对 FP4 来说一个乘法器只需要处理 4 位尾数实际上尾数只有 1 位芯片上可以放更小尺寸的乘法阵列。假设上一个架构用 4 位乘法器做 FP16需要 16 位乘法器执行两次 4 位乘法不合适的方式是FP4 位宽减半单次乘法需要的晶体管数量更少因此相同面积下能塞入更多并行单元。Rubin 大概率会针对 FP4 增加独立的张量核心通路或者让同一 Tensor Core 在 FP4 模式下按更大的 K 维度展开计算。这里有一个合理推断当前的 Tensor Core 指令通常以 m16n8k8 之类的 tile 为单位其中 k 是缩减维度。对于 16x8x8 的运算如果每个数都是 4 位数据位宽大约是 FP16 的四分之一那么内部计算阵列有能力把 k 扩展到 32 甚至 64一次性算完更长的规约链。这会让 FP4 GEMM 的指令级并行度进一步提升但也会给寄存器堆和累加器带来压力。2.3 带宽、显存与 PCBFP4 GEMM 的运行环境讲 FP4 GEMM 时很多人只盯着“算得快”但真实瓶颈经常是数据搬运。假设某个线性层的权重矩阵是 10000x10000FP16 下需要约 200MB 数据FP4 下只需要约 50MB但硬件吞吐翻倍后单位时间内消耗的数据量可能是原来的两倍多。也就是说FP4 不仅没有轻松反而更依赖显存带宽。所以 Rubin 必须搭配 HBM4、更宽的缓存和更高效的片上互连。热搜词里有“英伟达 rubin pcb”这也是非常关键的一点。PCB 和基板的设计直接影响信号完整性、电源完整性和散热能力。Chiplet 之间通过基板上的高速走线进行通信如果布线长度不匹配、串扰控制不好信号频率就提不上去。电源网络如果去耦电容不够瞬间电流波动会让 GPU 降频。换句话说FP4 的原始算力再高供电送不进去或者信号乱了实际表现也出不来。PCB 不是把芯片装起来那么简单它本质上是一个高速信号传输系统。Rubin 这类 2000W 级别功耗的 GPUPCB 上的电源层、地层设计、过孔位置、阻抗匹配都必须按许多 GHz 甚至几十 GHz 的标准去做。这是从 Blackwell 到 Rubin 一直都在强攻的硬骨头只是普通用户很难直接感知罢了。3. FP4 GEMM 的硬件实现思路3.1 FP4 数值格式为什么必须引入缩放因子FP4 的 E2M1 格式只有 1 位尾数符号 1 位指数 2 位指数偏置为 1。它可以表示的有限值数量非常少最大有限值约等于 6.0最小正常值约 0.5。如果待计算的数据分布在 0.001 到 1000 这个范围直接存成 FP4 几乎一半数据会变成 0 或 6.0损失惨重。解决方式是给数据加一个缩放因子。常见做法是按照 block 为单位比如 32 个元素一组为这一组元素计算一个公共的缩放因子然后用 FP4 存储规约后的数值。这个思想来自 OCP 制定的 Microscaling微缩放MX格式典型的有 MXFP4、MXFP8。缩放因子通常用 E8M0 这类纯指数格式存储只记录 2 的幂次这样硬件只需要做指数加减成本可以压得非常低。在 GEMM 当中A 矩阵和 B 矩阵各自有缩放因子相乘之后整体结果还要乘以两个缩放因子的乘积。因此 Tensor Core 的管线里不只是简单乘加而是要在累加结果上执行一个“缩放重加载”操作。这个操作可以放在 epilogue 阶段也可以提前在数据进入 Tensor Core 之前处理。无论哪种方式缩放因子的优化设计都决定了 FP4 GEMM 能带来多少收益。3.2 GEMM 流水线从加载到累加的完整路径一个 FP4 GEMM 在 GPU 上的执行过程大致可以分成以下几步全局内存读取 A、B 和各自的缩放因子写入共享内存从共享内存按 tile 读取经寄存器缓存后压入 Tensor CoreTensor Core 执行矩阵乘和累加得到 FP32 中间结果epilogue 阶段把 FP32 结果乘以缩放因子、加 bias、过激活函数最后按需转成 FP4 输出。与传统 FP16 GEMM 最大的不同在于FP4 经过的每个数据通路都要“窄一点”。例如从显存到共享内存的传输中一个 32 字节的请求可能携带 64 个 4 位元素。在寄存器阶段打包和解包 FP4 数据会成为额外开销。如果硬件不能高效处理这种半字节对齐软件层就需要自己做位压缩反而拖慢性能。因此 Rubin 的 Tensor Core 很可能在 Load 阶段就支持从内存读取紧凑的 4 位数据并在寄存器中自动解包成内部更宽的表示。另一种思路是把 FP4 直接映射到可寻址的 FP8 或 FP16 通道中比如一个 FP4 元素占半个 FP8 槽两个 FP4 拼成一个 FP8。这样硬件可以复用已有的低精度数据通路通过多路选择器和移位寄存器把两个 4 位值并行送入乘法器。这类实现属于微架构细节我们看不到但从软件性能参数上可以推测Rubin 的 FP4 峰值算力至少是 FP8 的两倍大概就是走了这类“位宽复用”的路线。3.3 2:4 稀疏 FP4稀疏 GEMM 的乘法吞吐NVIDIA 从 Ampere 架构开始支持 2:4 结构性稀疏每 4 个元素中只保留 2 个非零值且零值的位置固定。稀疏 GEMM 可以跳过一半的乘法运算理论上吞吐能翻倍。对 FP4 来说这种稀疏性同样重要因为 4 位权重本来就很小剪掉一半的权重大幅减少计算量。但稀疏性也有代价。稀疏矩阵必须额外存储索引或位掩码metadata在加载时需要读取额外数据。如果带宽本来就捉襟见肘稀疏化省下的计算时间可能被额外 metadata 的带宽开销抵消。实际中2:4 稀疏更适合权重矩阵因为权重可以提前离线剪枝而对于激活值由于每个 batch 不同很难固定零值位置所以很少用结构化稀疏。Rubin 大概率会延续 Blackwell 上的稀疏 Tensor Core 方案同时支持 FP4 和稀疏性混用。官方规格里的“峰值稀疏 FP4 TFLOPS”通常比稠密 FP4 高一倍但这个数字是理论极限真实场景能不能跑到一半以上要看数据分布和软件是否能把稀疏 pattern 高效映射到硬件。3.4 混合同精度 GEMM权重和激活精度不对称在实际部署中权重和激活值的敏感度往往不一样。有些层权重很稳定压到 FP4 几乎没有问题但激活值波动大用 FP4 会掉点需要保留 FP8。那么硬件能不能让 A 矩阵用 FP4B 矩阵用 FP8理论上 GEMM 的硬件乘法器应该能支持这种混合输入只需要把低精度的数据提升到同一个中间精度然后相乘。Hopper 的 FP8 模式就允许 A 和 B 使用不同的 E4M3/E5M2 格式。Blackwell 在 FP4 上很可能会扩展出混合精度组合比如 A:FP4, B:FP8或者 A:FP8, B:FP4。这样一来软件栈就可以按层选择最优的量化方案而不是一刀切全用 FP4。Rubin 的指令集如果能提供这种灵活性对精度和性能的平衡会带来很大帮助。不过这也意味着 Tensor Core 的内部乘法器需要支持不同位宽的输入通道硬件复杂度会上升。4. 软件栈与量化系统的配合4.1 从一行代码到 Tensor Core要经过多少层用户写一句torch.mm(a, b)看起来很简单但底层会经历 PyTorch 算子分派、ATen 矩阵乘选择、cuBLAS/cuBLASLt 库调用、CUDA 驱动最后才能落在 SASS 指令上。FP4 GEMM 的调用链比普通 FP16 更长因为中间还要传递量化参数、缩放因子和存储格式。如果应用程序直接调用 cuBLASLt需要配置一个matmulDescriptor指定数据类型为 FP4并提供缩放因子的维度信息。框架层的 TensorRT 或 vLLM 也都在为这种新精度做适配。就我目前看到的生态进展FP4 推理还没有达到 FP8 那样的成熟度很多量化工具都还在 preview 阶段。Rubin 真正发布后CUDA 工具包、cuBLAS、CUTLASS 都会快速跟进但在早期阶段工程师可能得自己写一些自定义 kernel 或者 workaround。关键建议是不要等硬件到了才开始调软件。现在就可以在模拟器或现有硬件上做 FP4 数值仿真确定哪些层适合 FP4、哪些层需要保留 FP8。软件栈成熟后迁移成本会低很多。4.2 缩放因子放哪里架构层面的 epilogue 设计FP4 GEMM 的输出节点通常要考虑 scale 和 bias 的融合。比如输出 C (A_fp4 * B_fp4) * (scaleA * scaleB) bias这个计算如果在 GEMM 内部完成可以避免多次访问显存。cuBLASLt 的 epilogue 设计支持不同组合不融、加 bias、加 scale、加激活函数等。在实际工程里我建议尽量把缩放因子的“数据类型”设计成 E8M0 这种纯指数整数因为它和任何尾数无关乘法就是指数相加非常简单。如果缩放因子也用 FP8 的 E4M3 格式硬件还得处理尾数乘法反而会引入额外误差和延迟。这不是一定的但 OCP MX 规范里 E8M0 就是设计来干这件事的。另一个细节是缩放因子的粒度。按 tensor 缩放最简单但精度最差按 channel也称为 per-channel压缩权重的精度更好按 block 缩放最强但需要额外存储。PyTorch 中常见的 weight_quant 实现就支持 per-group scale。在 FP4 GEMM 里缩放因子本身也是要参与计算的所以它的存储布局会影响 GEMM kernel 的读取效率。最理想的情况是缩放因子随矩阵一起按 tile 加载避免二次访存。4.3 QAT 训练模拟让 FP4 模型不掉点如果你只是把训练好的 FP16 模型直接转成 FP4大概率会掉点尤其是激活值比较大的任务。最好的做法是使用量化感知训练QAT在训练前向过程中用fake_quantize模拟 FP4 的舍入误差让模型自己去适应。一个常见误区是QAT 模拟时直接把 FP4 的结果当浮点用却忽略了缩放因子的动态更新。你需要用直通估计器STE让梯度绕过量化节点并且缩放因子要根据当前激活的统计量实时调整。像quantization.observer里的 MinMaxObserver、PerChannelMinMaxObserver 等等都可以用来估计缩放。但要注意FP4 的动态范围太窄单纯统计 min/max 可能会被离群值带偏。更稳妥的做法是使用百分位剪裁比如 99.9% 分位点作为缩放上限这样既保留大部分数据又容忍极端值溢出。QAT 训练中如果硬要精确模拟 FP4 GEMM 的硬件行为比较麻烦因为 Tensor Core 内部的舍入可能和软件模拟不一样。我通常建议先按 E2M1缩放因子的公式写一个并行 CPU/GPU reference再和预期结果对照。这个参考实现的性能不需要多好但正确性必须高。4.4 伪代码如何配置一个 FP4 GEMM 调用假设我们要在 cuBLASLt 里执行一个 FP4 GEMM伪代码逻辑如下// 这里以 cuBLASLt 风格示意实际 API 以最新版为准 cublasLtMatmulDesc_t opDesc; cublasLtMatmulDescCreate(opDesc, CUDA_R_32F, CUDA_R_32F); // 设置 A、B 均为 FP4输出为 FP32 cublasLtMatmulDescSetAttribute(opDesc, CUBLASLT_MATMUL_DESC_A_TYPE, CUDA_R_4F, sizeof(CUDA_R_4F)); cublasLtMatmulDescSetAttribute(opDesc, CUBLASLT_MATMUL_DESC_B_TYPE, CUDA_R_4F, sizeof(CUDA_R_4F)); // 指向缩放因子 cublasLtMatrixLayout_t scaleLayout; // 设置 scale 的维度、stride指向 scaleA 和 scaleB关键点是scaleLayout必须和矩阵的布局对齐。如果缩放是 per-channel 的那么 scaleA 的大小就是 M 或 K 的长度如果 per-block需要额外传入 block 尺寸。API 调对了底层 kernel 才能在所有 tile 上高效读取 scale。真实项目里我反而建议先用 TensorRT 的 Q/DQ 节点或 vLLM 的 FP4 实现进一步向下调封装那会让开发效率高很多。自己写 cuBLASLt 调用虽然灵活但要处理太多格式细节比如 FP4 数据的位序、端序、对齐稍不留神就出 bug。5. 工程实践中的性能调优与避坑5.1 先判断是带宽受限还是算力受限每次做 FP4 GEMM 优化我第一件事就是算计算强度arithmetic intensity。计算公式是计算强度 FLOPs / Bytes。如果这个值小于机器的“机械特性强度”那这个 GEMM 就是带宽受限反之则是算力受限。举个例子一个 4096 x 4096 的矩阵乘FLOPs 约等于2*4096^3 1.37e11。如果 A 和 B 都是 FP4每个元素只有 0.5 字节总共需要读取2*4096*4096*0.5 16.8MB。那么计算强度大约是 8155 FLOPs/byte。再看机器算力假设峰值 FP4 是 50 TFLOPS带宽是 8 TB/s那么机械特性强度是 6250 FLOPs/byte。这个例子中计算强度略高于机械特性它在理论上偏向算力受限。但一旦你引入 padding、metadata、非对齐访存实际有效带宽会打折很有可能变成带宽受限。所以优化方向在不同卡上完全不一样。带宽受限时优先做数据压缩、减少 metadata、使用 TMA 异步复制算力受限时优先调整 tile 大小、提高 Tensor Core 利用率、减少寄存器浪费。不能一把梭。5.2 半字节存储和内存对齐的麻烦FP4 数据是 4 位两个 FP4 放在同一个字节里。如果你按常规的uint8_t数组存储低位是第一个元素还是高位是第一个元素在不同 API 里可能有不同约定。CUTLASS 通常建议用uint8_t表示两个元素但具体 bit 顺序需要自己核对。更麻烦的是矩阵的 K 维或 M 维如果不对齐到 2 的倍数最后一个字节会有一半是空数据。你必须在分配内存时做 padding也就是把 stride 适当加大。我习惯把所有维度都对齐到 16 或 32 的倍数虽然会浪费一点显存但能避免很多 kernel 边界问题。另外一个坑是共享内存的 bank conflict。FP4 数据在共享内存中打包得非常紧凑读取时如果多个线程访问同一个字节的不同半字节可能会产生 bank conflict。最佳实践是把打包后的数据先按 uint8_t 数组加载到寄存器再在寄存器内部做位移和掩码而不是让共享内存硬件去猜你要哪个半字节。这样可控性最高性能也更稳定。5.3 稀疏性不是白拿的metadata 也会占带宽2:4 结构稀疏的 FP4 GEMM理论上乘法任务少一半但代价是要额外读取 metadata。每个 4 元素组用 4 位或 8 位保存哪些位置是非零值所以稀疏矩阵的存储从原来的 2 nibbles 变成了额外 metadata。如果稀疏矩阵本身非零值比例高或者零值分布不是 2:4 结构那就要重新压缩否则收益基本为零。我的建议是在模型权重上做一次全局稀疏率统计如果某层剪枝后只有 30% 的零值那硬套 2:4 会打乱数据带来精度损失但性能收益不明显。对于权重稀疏率超过 60% 的层再开 2:4 稀疏。对于激活值或 KV Cache除非专门做过结构化剪枝否则别碰稀疏。5.4 功耗、散热与供电对峰值性能的影响FP4 GEMM 的高吞吐会让 GPU 瞬时功耗明显上升尤其是跑那种大型密集矩阵乘时芯片内部数百个 Tensor Core 同时翻转电流冲击非常剧烈。Rubin 如果 PCB 上的电源稳压模块、去耦电容、陶瓷电容布局不合理电压纹波会比较大GPU 只能靠降频来保护硬件。我在之前做 Blackwell 性能测试时把功耗墙从 100% 调到 80%FP4 峰值算力骤降了 15% 以上而且 FP8 的下降幅度没那么明显。原因是 FP4 GEMM 的并行度更高、功耗波动更剧烈供电跟不上时 Boost 频率很难维持。因此评估 Rubin 的 FP4 性能一定要看持续功耗下的实际吞吐不要只看规格表峰值。如果你有整机柜部署计划电源和散热预算也要按 FP4 满负荷去设计留足冗余。6. FP4 GEMM 的场景落地与下一步6.1 场景图谱哪些任务真正需要 FP4我做了一个简单的适用性表格可以参考场景FP4 GEMM 适用性原因LLM 推理大量并行请求高权重和 KV Cache 都适合低精度压缩吞吐收益大图像生成Stable Diffusion 类中高UNet 和 Transformer 块的线性层很多但激活值波动大可能需要混合精度推荐系统大规模 embedding MLP高embedding 严重带宽受限FP4 能大幅减少内存占用科学计算低精度需求高FP4 太少自动驾驶边缘推理低延迟敏感FP4 对异常天气场景的鲁棒性不足这里有一个规律凡是“数据量大、算力相对充裕、精度容忍度较高”的任务都适合 FP4 GEMM。典型的就是大模型推理和推荐系统它们的高吞吐靠的是大批量把内存里的数据灌给计算单元而不是在这几个比特里抠精度。6.2 LLM 推理的数据流建议Prefill 用 FP8、Decode 用 FP4在 LLM 推理中Prefill 阶段要处理很长的 prompt计算强度高适合用 FP8 甚至 FP16 保持较高精度。Decode 阶段每次只生成一个 token主要从显存捞权重和 KV Cache带宽占用远大于计算需求这时候用 FP4 量化权重和 KV Cache 能显著降低延迟。Rubin 的软件栈如果要动态在 FP8 和 FP4 之间切换那会非常有意思。量化切换本身有开销但如果模型已经按层做了 QAT提前确定好哪些层用 FP4、哪些层用 FP8运行时就不需要重新量化只需要切换 GEMM kernel 即可。这也是批量推理引擎将来最需要的功能之一。6.3 对未来训练和微调的影响FP4 GEMM 直接用稳定训练目前还比较困难。反向传播需要计算梯度和更新权重梯度范数变化很大4 位浮点的动态范围不够。不过有一些研究方向在尝试“低精度梯度压缩”用 FP4 近似梯度再配合误差反馈补偿这样也许能在分布式训练中减少通信量。微调场景下可以采用“冻结权重为 FP4 表示但保留一份 FP32 的影子权重”的做法。前向推理用 FP4 GEMM反向传播时用影子权重的高精度梯度。这本质上是一种混合精度微调显存占用虽然增加了但能让大模型在有限的 GPU 上做轻量适配。Rubin 的高带宽架构对这种“影子权重”模式很友好因为后面更新梯度需要读取 FP32 权重带宽如果不够反而会拖慢训练。结尾一点个人体会最后说点个人经验。我在做量化部署时第一眼看到 FP4 觉得真是好东西但实际踩坑之后发现FP4 GEMM 的成功不是靠硬件一个维度而是数值格式、软件栈、供电和散热整体配合的结果。我的建议是在 Rubin 真正大规模出货之前先把你的模型在 FP4 模拟器上跑一遍记录误差曲线和内存分布。不要等硬件到了才开始调内核那时候就晚了。另外一个实用技巧是如果你打算用 FP4 GEMM尽量把缩放因子设计成 8 位指数E8M0这样硬件可以用极低成本完成缩放还能减少数值溢出的概率。后面我也会持续跟进 Rubin 的相关资料有新发现再和大家分享。