【免费下载链接】DeepSeek-V4-Pro-0813-nvfp4-DSpark项目地址https://ai.gitcode.com/hf_mirrors/nvidia/DeepSeek-V4-Pro-0813-nvfp4-DSpark点击查看免费下载DeepSeek-V4-Pro-0813-nvfp4-DSpark是 DeepSeek-V4-Pro-0813 的 NVFP4 量化版本而它推理代码的灵魂就藏在 inference/kernel.py 这一个文件里稀疏注意力内核、Sinkhorn 归一化内核、FP4 GEMM 内核全部基于 TileLang 手写。本文带你逐个读懂这些高性能算子看看一个 1.65T 参数的 MoE 模型是如何用极小的显存与极高的吞吐跑起来的。kernel.py 全景一个文件承载全部核心算子 整个推理目录非常精简inference/ 下只有 4 个 Python 文件。其中 kernel.py 定义了5 组 TileLang 内核覆盖了量化、矩阵乘、注意力的完整链路内核函数位置作用act_quant_kernelL40-L102每 128 元素分块的 FP8 激活量化fp4_quant_kernelL128-L183每 32 元素分块的 FP4 量化E8M0 缩放fp8_gemm_kernelL203-L254FP8×FP8 分块缩放矩阵乘sparse_attn_kernelL276-L352top-k 稀疏注意力 在线 softmaxhc_split_sinkhorn_kernelL371-L427超连接权重拆分 Sinkhorn 迭代fp4_gemm_kernelL441-L515FP8 激活 × FP4 权重混合精度矩阵乘所有内核共用一组编译开关 L9-L12关闭 warp specialization 与 TMA lowering换取更稳定的编译行为。文件开头还定义了四个关键数据类型别名——FP8、FP4、FE8M0、BF16L14-L19它们是理解 NVFP4 体系的钥匙。快速上手环境依赖与运行方式仓库是只读的获取源码git clone https://gitcode.com/hf_mirrors/nvidia/DeepSeek-V4-Pro-0813-nvfp4-DSpark依赖很简单见 inference/requirements.txttorch2.10.0、tilelang0.1.8版本被严格锁定因为 TileLang 编译器 API 变化较快以及fast_hadamard_transform。按 inference/README.md 的流程先用 convert.py 把权重切分到多卡再用 generate.py 启动对话。关键点在 inference/config.json 里的expert_dtype: fp4L37——它决定 MoE 专家权重以 FP4 格式加载此时每次线性层都会走到本文要讲的fp4_gemm内核。量化内核FP8 与 FP4 如何块级压缩数据FP8 分块量化amax 归一 可选 2 的幂缩放act_quant_kernelL40-L102把激活值按block_size128分块每块先算绝对值最大值T.reduce_absmax用amax / 448得到缩放因子把数据压进 FP8 的 ±448 范围。它有两个很实用的开关round_scaleTrue把缩放因子向上取整为 2 的幂得到 MXFP 风格的E8M0缩放只需 8 位存储。这个取整不靠慢速的log/ceil而是用 fast_log2_ceil 直接操纵 IEEE 754 位模式——取出指数字段、检查尾数是否非零一行位运算搞定inplaceTrue量化后立刻反量化回 BF16用于模拟量化噪声QAT 风格model.py 里对 KV 的量化模拟正是这样做的L512。FP4 分块量化32 元素一块E8M0 缩放fp4_quant_kernelL128-L183是 NVFP4 权数量化逻辑的参照实现每 32 个元素共享一个E8M0缩放数据值域限在 ±6.0E2M1 的最大值。两个细节值得注意amax有6 * 2^-126的下界保护L163防止全零块除零缩放同样走 fast_round_scale 的 2 的幂取整。FP8 GEMM双累加器提升有效精度fp8_gemm_kernelL203-L254计算C A B^TA、B 都是 FP8各自带每 128 元素的块级缩放。它的工程亮点有三个流水线深度 4T.Pipelined(K_iters, num_stages4)让数据搬运、FP4/FP8 转换与T.gemm计算重叠L2 缓存优化T.use_swizzle(panel_size10)L233打乱块执行顺序提高共享缓存命中双累加器C_local做原始 FP8 点积C_local_accum单独累加乘过缩放因子的结果L248-L249注释直接写明这是 2x accumulation precision——把缩放修正放在独立累加器里避免每步都混入低精度数据。FP4 GEMMFP8 激活 × FP4 权重的混合精度矩阵乘 ⚡fp4_gemm_kernelL441-L515是 NVFP4 检查点的核心路径激活是FP8每 128 一个缩放权重是FP4每 32 一个 E8M0 缩放。它的策略非常聪明L453-L456 的注释说明了全部思路借用 FP8 张量核心GPU 的矩阵乘单元对 FP8 有原生加速FP4 没有单独的 GEMM 路径。于是内核把 FP4 子块加载进 shared memory 后先经 FP32 再转成 FP8L494-L496注释解释绕开 C 重载二义性然后执行一次标准 FP8×FP8 GEMMblock_K32与权重缩放对齐每个 K 块恰好是一个权重缩放组scales_b按k直接索引L498-L500激活缩放组是 128所以scales_a按k // 4索引L502-L504每 4 个子块复用同一个激活缩放缩放修正融入累加C_local_accum C_local * scale_a * scale_bL508-L509与 FP8 GEMM 的双累加器技巧一脉相承。Python 侧入口 fp4_gemm 还说明了权重在内存中的布局[N, K//2]的float4_e2m1fn_x2打包格式2 个 FP4 值挤进 1 个字节。稀疏注意力TileLang 版 FlashAttention 索引收集 sparse_attn_kernelL276-L352是这个模型混合注意力Compressed Sparse Attention Heavily Compressed Attention的执行引擎。它接收q、kv和一张top-k 索引表topk_idxs每个序列位置只计算索引命中的 KV把长上下文的注意力复杂度从 O(n²) 压到 O(n·topk)。内核内部是标准的 FlashAttention 式在线 softmax配合索引收集按索引 gather每个 64 元素的 KV 块先取出对应索引越界补 -1无效位置填 0、分数填-infL321-L327两次 GEMMQ K^T得分数再P V累加输出中间全程在寄存器/共享内存里完成无需写出完整注意力矩阵在线数值稳定滚动维护scores_max与sum_exp每来一个新块就用exp(旧max - 新max)修正历史累积L331-L342可学习的attn_sink结束时把exp(attn_sink - max)加进分母L345-L346等价于给注意力分布加一个逃生通道避免 softmax 被迫在无关 KV 上摊薄权重。它在模型里的用法见 Attention.forward滑动窗口window_size128的 KV 恒被选中压缩 KV 的 top-k 索引由 Indexer 学出来index_topk1024两者拼接后交给sparse_attnL513-L538。DSpark 投机解码头也复用同一个内核L785。Sinkhorn 内核让超连接权重行行归一、列列归一hc_split_sinkhorn_kernelL371-L427服务于模型的另一大特色——超连接Hyper-Connections隐藏状态不是一条残差流而是hc_mult4份拷贝并行流动每层要用学到的权重在多份 → 1 份 → 多份之间混合。每个 block 从(2hc)*hc 24个混合参数里拆出三种权重presigmoid 门控 偏置把 4 份残差加权压成 1 份L391-L392post2·sigmoid门控把 1 份输出放大回 4 份L393-L394comb4×4 的组合矩阵负责 4 份残差之间的相互混合L395-L396。comb的归一化是内核的精华先按行做数值稳定的 softmax减行最大值再 expL401-L408然后进入Sinkhorn 迭代——交替执行除以行和除以列和L415-L423。sinkhorn_iters20config.json L34轮迭代后矩阵的行和与列和都趋近于 1得到一个双随机的混合矩阵——这正是 DeepSeek 论文里 Manifold-Constrained Hyper-Connections 的约束保证残差流动在训练中长期稳定。整个 4×4 矩阵常驻寄存器 fragmentcomb_frag20 轮迭代零显存往返单线程块即可完成。Python 侧 hc_split_sinkhorn 被 Block.hc_pre 调用每个 Transformer 块的 attention 与 FFN 各用一次。内核与模型如何协作一张调用链全景 ️把 model.py 与 kernel.py 合起来看调用链非常清晰模型组件调用的内核说明linear()分发器L114-L126act_quantfp4_gemm/fp8_gemm按权重 dtype 自动选择精度路径AttentionL442-L548sparse_attn窗口 KV top-k 压缩 KV 的稀疏注意力BlockL652-L716hc_split_sinkhorn每次 attention/FFN 前后的残差混合DSparkAttentionL750-L792sparse_attn投机解码的块状草稿注意力想改参数全部在 inference/config.jsonhc_mult、hc_sinkhorn_iters、window_size、index_topk、expert_dtype都是明文 JSON 键ModelArgs 逐字段对应。总结这份 kernel.py 值得抄的作业位运算替代数学函数fast_log2_ceil/fast_pow2用 IEEE 754 位模式完成 2 的幂取整是低精度量化内核的通用技巧L22-L37没有原生 FP4 GEMM就用 FP8 张量核心FP4→FP8 上转换 分块缩放对齐思路可迁移到任何新精度 老硬件的场景在线 softmax 索引 gather把 FlashAttention 的滚动归约与稀疏索引收集合进一个内核长上下文注意力不再受制于完整 KV小矩阵常驻寄存器4×4 的 Sinkhorn 迭代完全在 fragment 里完成 20 轮是小计算量、高迭代次数算子的最佳实践。从 66 个 safetensors 分片到 model-dspark-mtp.safetensors 的草稿头这个 NVFP4 检查点能高效运行的秘密一半在量化配方另一半就在这 500 多行 TileLang 内核里。读懂它你也就读懂了当前超大 MoE 模型推理优化的主流打法。赞分享【免费下载链接】DeepSeek-V4-Pro-0813-nvfp4-DSpark项目地址https://ai.gitcode.com/hf_mirrors/nvidia/DeepSeek-V4-Pro-0813-nvfp4-DSpark点击查看免费下载相关推荐4比特权重如何计算DeepSeek-V4-Flash-nvfp4-DSpark的TileLang FP4×FP8混合GEMM与量化内核深度剖析4比特权重如何计算DeepSeek V4 Flash nvfp4 DSpark的TileLang FP4×FP8混合GEMM与量化内核深度剖析 DeepSeeTileLang 实现 DeepSeek V3.2 稀疏 MLA 全流程Lightning Indexer、Top-k Selector 与稀疏注意力内核实战TileLang 实现 DeepSeek V3.2 稀疏 MLA 全流程Lightning Indexer、Top k Selector 与稀疏注意力内核实战编程语言编译器高性能计算人工智能1M 上下文如何炼成深度拆解 DeepSeek-V4-Pro-0813-nvfp4-DSpark 的混合注意力与 YaRN 长文本机制1M 上下文如何炼成深度拆解 DeepSeek V4 Pro 0813 nvfp4 DSpark 的混合注意力与 YaRN 长文本机制 本文以 DeepSee上一篇Viper 删除主机完全指南主机与 Session 的生命周期管理下一篇鸣潮工具箱 WaveTools一条命令装好120帧、画质与抽卡保底一次管明白创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考