Repository navigation
Conversation
rtp-llm-review-bot
added
auto-lgtm:needs-approval
自动审查进度
更新时间：2026-10-23 18:53 北京时间。运行期间约每 5 分钟更新；阶段任务数会随核验展开，不能当作整轮完成百分比。 |
rtp-llm-review-bot
added
review:delayed
rtp-llm-review-bot
added
review:running
rtp-llm-review-bot
added
review:delayed
| # experts under ``layers.*``. Keep both layouts available and let the | ||
| # runtime quantization setting decide which kernels consume them. | ||
| self._raw_hf_stacked = any( | ||
| ".mlp.experts.gate_up_proj" in key for key in weight_keys |
There was a problem hiding this comment.
[P1] raw HF 堆叠路径使 DSV4 路由专家落入无 swiglu_limit 钳制的 FP8_PER_BLOCK DeepGemm 执行器族,激活语义静默偏离 V4 契约
加载含 .mlp.experts.gate_up_proj 的微调 raw HF DSV4 导出(config.json 携带 swiglu_limit>0,即架构文档所述 V4 语义;作者对本路径 runtime config 的新测试取 7.0 且 moe_quant_method="FP8_PER_BLOCK"、moe_strategy="fp8_per_block_pure_dp"),BF16 权重触发报错后按提示配置 --quantization FP8_PER_BLOCK,在 SM90+ 且 deep_gemm 可用环境以 auto 或任一 fp8_per_block_* 策略服务任意请求;启用本 PR 的 norm Top-p 特性时守卫强制落入该执行器族。
base 中 Fp8Fp4MoeRuntimeConfig.moe_quant_method 为固定类属性 "FP8_FP4"(base layer.py:27,构造器 :35-57 无该参数)且 quant_config=None(:113),config_resolver.get_quant_method 恒返 None(:62-63,该文件未变更两版相同),base fp8_per_block 策略要求 FP8_PER_BLOCK(base fp8_per_block.py:30-31)故对 DSV4 双重不可达;base deepseek_v4.py 仅构建 FP4 路由描述符(:495),全文件无 gate_up_proj/_raw_hf_stacked(检索零命中)。DeepGemm 执行器无钳制属既有代码,但 head 新增堆叠描述符与检测分支(deepseek_v4.py:159-161/:505-542/:594-597)、dtype 探测与引导报错(layer.py:230-240)、quant_config 构造(:140-144)、norm Top-p 守卫白名单仅放行 fp8_per_block 族(:271-283)、attention int32 标度适配(:377-391),首次把 DSV4 路由专家导入该无钳制执行器族,且服务接线完整(:788-800 构造 DeepSeekV4Model,文件头 M0 注释为过期信息)。
head 的 _process_meta 检测到 .mlp.experts.gate_up_proj(deepseek_v4.py:160)即走 _build_routed_experts_stacked(:505-542,:594-597)产出 BF16 W.moe_w1/w2;Fp8Fp4MoeLayer 对 BF16 报错并引导 --quantization FP8_PER_BLOCK(layer.py:236-240);该 CLI 值经 quant_config.py:1086-1095 回退取 DEFAULT_FP8_BLOCK_WISE_QUANT_CONFIG(is_quanted=False,:1039-1043/:1071),is_quanted=True 的 PerBlockFp8Weight/V4PerBlockFp8Weight(:323-326/:1188-1191)不竞争,W.moe_w1(在 w8a8_weight_list,:310)唯一命中 LoadQuantPerBlockFp8Weight(:913-918),经 per_block_cast_to_fp8 在线量化为 float8_e4m3fn(:169,:1086);layer 依 dtype 注册 moe_quant_method=FP8_PER_BLOCK(:230-232,:264)并构造 Fp8BlockWiseQuantConfig(is_quanted=True)(:140-144),使 resolver.get_quant_method 返回 FP8_PER_BLOCK(quant_config.py:549-550,config_resolver.py:62-64)。注册表先按 supported_moe_quant_method 预过滤(strategy_registry.py:18-28,:99-114)再经 can_handle 三重检查(strategy_base.py:89-95):FP8_FP4 族(fp8_fp4.py:18/:30,mega_moe/grouped_fp4/local_loop,均带钳制)被排除,batched_triton 要求无量化(:43),cutlass 要求 per-tensor 方法(:53-55),pure_tp per-tensor 路由要求 per-tensor 方法(:198-201),故仅六个 fp8_per_block 策略存活(fp8_per_block.py:31/:67/:112/:154/:192/:226)并绑定 DeepGemmHybrid/MaskedV2/Masked(:51/:94/:135/:178/:212/:246;V2 继承 hybrid 并复用 execute_masked,v2:22/:60)。该执行器族全部激活路径(deepgemm_hybrid_executor.py:426 triton、:588/:607 masked、:850 contiguous;masked 执行器 :235/:254)调用 common/activation.py 的裸 silu(gate)*value 内核(:34-40,函数签名 :43-45/:574-580/:811-817/:1045-1054 均无钳制参数,文件内 5 处 tl.clamp 全为 FP8 量化范围钳制 :139/:410/:560/:740/:1031;dsv41 私有内核注释明言框架内核缺 swiglu_limit 钳制参数而 V4 需要)。cfg.swiglu_limit 虽传入 runtime config(layer.py:248)且被同层共享专家消费(:306,shared_expert.py:165),路由专家调用仅传 activation="silu"(:381-394),钳制被静默丢弃。
每个 MoE 层每 token 的路由专家 FFN 激活缺失训练时的钳制非线性,与参考 Expert 及转换 FP4 路径(带钳制)产生系统性数值偏差,生成输出静默劣化且无报错无日志;同层共享专家仍钳制,路由/共享贡献相互不一致。属受支持非默认路径上的错误输出类运行时后果。
rtp_llm/models/deepseek_v4.py:160：".mlp.experts.gate_up_proj" in key for key in weight_keys
rtp_llm/models/deepseek_v4.py:595：weights += self._build_routed_experts_stacked(layer_id)
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:232：moe_quant_method = "FP8_PER_BLOCK"
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:239："--quantization FP8_PER_BLOCK."
rtp_llm/model_loader/per_block_fp8_quant_weight.py:918：return name in cls.w8a8_weight_list and name not in [W.mla_kc, W.mla_vc]
rtp_llm/model_loader/per_block_fp8_quant_weight.py:310：W.moe_w1: W.moe_s1,
rtp_llm/model_loader/per_block_fp8_quant_weight.py:1086：quant_kernel, scale = per_block_cast_to_fp8(
rtp_llm/model_loader/per_block_fp8_quant_weight.py:169：x_scaled = (x_view * (FP8_E4M3_MAX / x_amax)).to(torch.float8_e4m3fn)
rtp_llm/config/quant_config.py:1042：is_quanted=False,
rtp_llm/config/quant_config.py:1071："FP8_PER_BLOCK": DEFAULT_FP8_BLOCK_WISE_QUANT_CONFIG,
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:141：Fp8BlockWiseQuantConfig(is_quanted=True)
rtp_llm/config/quant_config.py:550：return "FP8_PER_BLOCK"
rtp_llm/models_py/modules/factory/fused_moe/strategy_registry.py:23：return supported == quant_method
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_fp4.py:30：checker.check(config.moe_quant_method == "FP8_FP4")
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_per_block.py:31：checker.check(quant_method == "FP8_PER_BLOCK")
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_per_block.py:51：executor_class=DeepGemmHybridExecutor,
rtp_llm/models_py/modules/factory/fused_moe/impl/common/executor/batched_triton_executor.py:43：checker.check(not resolver.has_quantization(config))
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/cutlass_moe.py:54：quant_method in ["FP8_PER_TENSOR_COMPRESSED", "FP8_DYNAMIC_PER_TENSOR"]
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_hybrid_executor.py:426：silu_and_mul(down_input, gate_up_output)
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_hybrid_executor.py:607：silu_mul_masked_fp8_post_quant_fwd(
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_hybrid_executor.py:850：silu_and_mul(down_input, gateup_output)
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_masked_executor_v2.py:60：return self.execute_masked(
rtp_llm/models_py/triton_kernels/common/activation.py:37：silu_gate = gate * tl.sigmoid(gate.to(tl.float32))
rtp_llm/models_py/modules/dsv41/moe/_silu_mul_fp8_quant_triton.py:20：the framework kernel has no swiglu_limit clamp parameter (V4 needs it)
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:248：swiglu_limit=swiglu_limit,
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:386：activation="silu",
rtp_llm/models/deepseek_v4.py:21： - SwiGLU clamping (linear in [-swiglu_limit,limit], gate <= limit)
rtp_llm/models/deepseek_v4.py:922：config.swiglu_limit = float(config_json.get("swiglu_limit", 0.0))
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/expert.py:130：clamp_limit=self.swiglu_limit,
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/local_loop.py:433：# SiLU + (optional clamp) + mul, in fp32 (matches Expert.forward).
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/grouped_fp4.py:274：clamp_limit=cfg.swiglu_limit,
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/shared_expert.py:165：clamp_limit=self.swiglu_limit,
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:306：swiglu_limit=swiglu_limit,
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:27：moe_quant_method = "FP8_FP4"
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:113：self.quant_config = None
rtp_llm/models_py/modules/factory/fused_moe/utils/config_resolver.py:63：return None
rtp_llm/models/deepseek_v4.py:495：weights += self._build_routed_experts_fp4(layer_id)
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:280："DeepSeek-V4 norm Top-p requires an FP8-per-block MoE strategy "
rtp_llm/models_py/modules/dsv4/transformer.py:88：swiglu_limit: float = 10.0
rtp_llm/models/test/moe_cuda13_config_propagation_test.py:471：swiglu_limit=7.0,
rtp_llm/models_py/modules/factory/fused_moe/utils/test/fp8_fp4_layer_test.py:215：(4, 16, 4), dtype=torch.float8_e4m3fn
rtp_llm/models_py/modules/factory/fused_moe/utils/test/fp8_fp4_layer_test.py:210：swiglu_limit=1.0,
rtp_llm/models/deepseek_v4.py:791：self.py_model = DeepSeekV4Model(
rtp_llm/models_py/model_desc/deepseek_v4_model.py:323：swiglu_limit=float(model_config.swiglu_limit),
rtp_llm/models_py/modules/dsv4/block.py:169：swiglu_limit=swiglu_limit,
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:207：self._moe = Fp8Fp4MoeLayer(
建议：为 FP8_PER_BLOCK DeepGemm 执行器族读取 cfg.swiglu_limit 并传入带钳制的激活内核(参照 dsv41 私有 _silu_mul_fp8_quant 端口,或为 common/activation.py silu_and_mul 系内核增加 clamp_limit 变体);短期可在 Fp8Fp4MoeLayer 中对 model_type=deepseek_v4 且 swiglu_limit>0 而所选执行器不支持钳制时显式报错;并补充 raw FP8_PER_BLOCK 路径与转换 FP4 路径的前向数值对照测试。
评审版本：f8c79ba72870
| moe_config=moe_config, | ||
| ) | ||
| args.parallelism_config = parallelism_config | ||
| args.moe_config = moe_config |
There was a problem hiding this comment.
[P2] 单卡 DeepSeek-V4（FP8-per-block routed 权重）在显式 USE_ALL_GATHER=0 下无任何可选 MoE 策略，模型构造期 get_strategy 抛 ValueError 致服务启动失败
单卡（tp=ep=dp=world_size=1）部署 DeepSeek-V4：raw HF stacked BF16 检查点配 --quantization FP8_PER_BLOCK（layer.py:236-239 报错文案明示的受支持形态），moe_strategy 默认 auto，启动参数或环境显式 USE_ALL_GATHER=0（受支持旗标，默认 True）。start_server.py:619 经 setup_and_configure_server→auto_configure_deepep 后 use_all_gather 保持 False、单卡 auto 分支 use_deepep_moe=False（server_config_setup.py:200），同一 moe_config 对象经 engine_config.py:187 进入模型构造。
base：dsv4 routed 权重仅来自 int8 打包 FP4（models/deepseek_v4.py:495），moe_quant_method 为类属性 'FP8_FP4'（layer.py:27），moe_config 为合成 SimpleNamespace 无 use_all_gather 字段（layer.py:119），PureTpRouterBase 的既有检查（base pure_tp_router.py:66）对 dsv4 不可达。head：新增 raw HF stacked BF16 装载分支（models/deepseek_v4.py:595）与 e4m3 在线量化（per_block_fp8_quant_weight.py:835），dtype 派生 'FP8_PER_BLOCK'（layer.py:232），并透传真实 moe_config（deepseek_v4_model.py:378、layer.py:155），使显式 0 值首次进入 dsv4 单卡策略判定。
head 将真实 moe_config 按引用写入 V4Args（deepseek_v4_model.py:378），经 transformer.py:190→block.py:178→chunked_layer.py:225 下发至 Fp8Fp4MoeRuntimeConfig（layer.py:155，非 None 时不走合成 fallback）；layer.py:230-232 按 routed 权重 dtype（e4m3）派生 moe_quant_method='FP8_PER_BLOCK'，model_config.quant_config 置为 Fp8BlockWiseQuantConfig（layer.py:141，get_method()='FP8_PER_BLOCK'，quant_config.py:550）。策略选择逐族核对：声明 'FP8_FP4' 的 mega/grouped_fp4/local_loop 族被 registry quant 过滤（strategy_registry.py:99-104、fp8_fp4.py:18）；can_handle 汇入 router 检查（strategy_base.py:93），fp8_per_block 族单卡可用的 NoDP/NoDPMasked 共用 PureTpRouterBase，use_all_gather=False 直接失败（pure_tp_router.py:66、config_resolver.py:151）；EpNormal/EpLowLatency 要求 ep_size>1（deepep_normal_router.py:48、deepep_low_latency_router.py:54、config_resolver.py:76）；PureDP/PureCP 要求显式 moe_strategy 与 dp/ep>1（fp8_per_block.py:118）；no-quant 要求 quant_method 为 None（no_quant.py:31）、batched-triton 拒绝量化（batched_data_router.py:50）、per-tensor/w4a8/fp4 因量化方法不匹配出局。候选集为空，get_strategy 抛 ValueError（strategy_registry.py:174），经 factory.py:53 从 Fp8Fp4MoeLayer 构造传播至 DeepSeekV4Model 初始化（models/deepseek_v4.py:791），启动失败。
运行时启动失败（非测试失败）：DeepSeekV4Model 构造期间 FusedMoeFactory().create_fused_moe→get_strategy 抛 ValueError 'No suitable MOE strategy found for configuration'，服务无法启动；报错文案只提示 quant_config/ep_size/parallelism，未点名 use_all_gather，排障被误导。int8 打包 FP4 检查点路径（layer.py:234 仍派生 'FP8_FP4'）与默认 USE_ALL_GATHER=True 部署不受影响；无静默错误输出或内存问题。定级 P2 而非 P1：该失败为启动期 fail-fast，触发需显式非默认旗标，且同一旗标组合在 base 的 generic per-block 模型上本已同样失败（base generic_moe.py:86 传真实 moe_config + base pure_tp_router.py:66 同一检查），本 PR 未新引入该契约违背类别，只是将新 dsv4 per-block 路径纳入既有契约。
rtp_llm/models_py/model_desc/deepseek_v4_model.py:378：args.moe_config = moe_config
rtp_llm/models_py/modules/dsv4/transformer.py:190：moe_config=args.moe_config,
rtp_llm/models_py/modules/dsv4/block.py:178：moe_config=moe_config,
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:225：moe_config=moe_config,
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:155：self.moe_config = moe_config or SimpleNamespace(
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:232：moe_quant_method = "FP8_PER_BLOCK"
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:141：Fp8BlockWiseQuantConfig(is_quanted=True)
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:161：use_all_gather=use_pure_dp or self.ep_size == 1,
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_tp_router.py:66：checker.check(resolver.use_all_gather(config))
rtp_llm/models_py/modules/factory/fused_moe/utils/config_resolver.py:151：return config.moe_config.use_all_gather if config.moe_config else True
rtp_llm/config/server_config_setup.py:107：moe_config.use_all_gather = (
rtp_llm/config/server_config_setup.py:200：use_deepep_moe = False
rtp_llm/server/server_args/moe_group_args.py:187：default=True,
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/deepep_normal_router.py:48：checker.check(resolver.is_ep_enabled(config))
rtp_llm/models_py/modules/factory/fused_moe/utils/config_resolver.py:76：return config.ep_size > 1
rtp_llm/models_py/modules/factory/fused_moe/strategy_registry.py:174："No suitable MOE strategy found for configuration. "
rtp_llm/models_py/modules/factory/fused_moe/strategy_registry.py:99：quant_method = getattr(config, "moe_quant_method", None)
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_fp4.py:18：supported_moe_quant_method = "FP8_FP4"
rtp_llm/models_py/modules/factory/fused_moe/defs/strategy_base.py:93：router_cls.check_conditions(checker, config)
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/no_quant.py:31：checker.check(quant_method is None)
rtp_llm/models_py/modules/factory/fused_moe/impl/common/router/batched_data_router.py:50：checker.check(not resolver.has_quantization(config))
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_per_block.py:118：checker.check(config.moe_strategy == "fp8_per_block_pure_dp")
rtp_llm/config/quant_config.py:550：return "FP8_PER_BLOCK"
rtp_llm/models/deepseek_v4.py:595：weights += self._build_routed_experts_stacked(layer_id)
rtp_llm/model_loader/per_block_fp8_quant_weight.py:835：data_type=torch.float8_e4m3fn,
rtp_llm/config/engine_config.py:187：moe_config = py_env_configs.moe_config
rtp_llm/start_server.py:619：setup_and_configure_server(py_env_configs)
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:27：moe_quant_method = "FP8_FP4"
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:119：self.moe_config = SimpleNamespace(
rtp_llm/models/deepseek_v4.py:495：weights += self._build_routed_experts_fp4(layer_id)
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_tp_router.py:66：checker.check(resolver.use_all_gather(config))
rtp_llm/models_py/model_desc/generic_moe.py:86：moe_config=moe_config,
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/fp8_fp4_router.py:29：checker.check(config.moe_quant_method == "FP8_FP4")
rtp_llm/models_py/modules/factory/fused_moe/factory.py:53：strategy = self.registry.get_strategy(config)
rtp_llm/models/deepseek_v4.py:791：self.py_model = DeepSeekV4Model(
建议：任选其一：1) Fp8Fp4MoeRuntimeConfig 对 ep_size==1 沿用 fallback 语义（use_all_gather=use_pure_dp or ep_size==1），不透传单卡下无通信意义的 0 值；2) auto_configure_deepep 对 is_single_gpu 恢复 use_all_gather=True 并记日志，与 server_config_setup.py:104-106 的守护意图一致；3) 放宽 PureTpRouterBase 在 ep==1 时对 use_all_gather 的硬性要求；4) 至少在 'No suitable MOE strategy' 报错文案中点名 USE_ALL_GATHER 与单卡拓扑冲突，避免误导排障。
评审版本：f8c79ba72870
| plan = torch.tensor( | ||
| [int(tokens), local_chunk_count], dtype=torch.int64, device=device | ||
| ) | ||
| dist.all_reduce(plan, op=dist.ReduceOp.MAX, group=dist.group.WORLD) |
There was a problem hiding this comment.
[P1] eager decode 前向新增 WORLD 计划 all_reduce 而 CUDA 图回放 rank 永不参与；已交付 mega_moe_se+enable_cuda_graph+dp2/ep2 解码配置中图/eager 按 rank 本地分歧时集合错位，整实例挂死或崩溃
仓库已交付的 decode 角色配置(base/head suites_cuda13.bzl:65 逐字相同)：『--enable_cuda_graph 1 --decode_capture_config '1,2,4,8,16' --tp_size 1 --dp_size 2 --ep_size 2 --moe_strategy mega_moe_se』；concurrency_limit 默认 32 且 max_generate_batch_size 由其赋值，DECODE 调度按该上限准入(FIFOScheduler.cc:145)，而固定捕获阶梯封顶 16。某解码步 rank A 本地批量 17..32 超出最大捕获桶，按 cuda_graph_runner.cc:1291-1299 本地回退 eager 并进入计划 all_reduce；rank B 批量命中桶(或为假流 bs=1)回放图。canReplaySelectedGraph 的其余本地拒绝条件(tag 不匹配、position 校验失败、key 未捕获)同样可制造分歧；head 空批路径(484)在对端回放时同样阻塞。
base needs_sync 仅限 {mega_moe, mega_moe_se} 且显式排除 is_decode_role(base chunked_layer.py:108-109)，synchronized_chunk_tokens 对 decode 角色/decode 前向提前返回本地 tokens(base:235-236)，decode 层循环为裸循环(base decode/forward.py:347)，空批直接返回零行输出(base:470)；base 中 synchronized_moe_chunk_plan 仅见于 transformer.py:683 与 prefill/forward.py:622。head 移除 decode-role 排除、策略集扩为四类(head:29-36)、计划改为 [tokens, chunk_count] 双元素 WORLD all_reduce(head:148-153)，并由 decode/forward.py:351/484 与 transformer.py:571 包裹进解码路径。
head 的 synchronized_moe_chunk_plan 在 needs_sync(策略∈_COLLECTIVE_CHUNK_STRATEGIES 且 chunking_enabled，已无 decode-role 排除)成立时，于层循环前执行 dist.all_reduce(plan, MAX, WORLD) 并以两次 .item() 阻塞；decode/forward.py 层循环(351)、空批路径(484)与 transformer.forward_decode(571)均包裹该计划。图捕获期 121 行守卫跳过归约(捕获在 cuda_graph_runner.cc:2041 执行 Python，计划从未入图)，回放路径 replayDecode→replayGraph→graph.replay() 只重放内核、不执行任何 Python 也无 host 集合，回放 rank 在该步永不发起此 all_reduce。canRun 完全按本 rank 输入判定：本地批量超出最大捕获桶即按设计回退 eager(cuda_graph_runner.cc:1291-1299)，DP 各 rank 批量天然不均、空闲 rank 由假流补 bs=1 照常进步；cpp/cuda_graph 全文无分布式调用，唯一的跨 rank 同步助手仅限 warmup(环境门控)。同一步 rank A 回退 eager 阻塞在计划 all_reduce(NCCL 超时被有意设为约百年不自愈)，rank B 回放图，其 mega 对等 NVLink 内核要求全员同时进入，等不到 rank A 即 NVLink barrier 超时触发陷阱；双方互等，整引擎挂死/崩溃。
rank A 的 WORLD all_reduce 永远凑不齐成员，主机线程阻塞在 .item()(NCCL 超时约百年、有意阻塞不自愈)；rank B 图内 mega 内核等不到 rank A，NVLink barrier 超时→SIGTRAP，表现为引擎级 CUDA_ERROR_LAUNCH_FAILED；整实例挂死/崩溃，需外部重启，属运行时故障而非测试失败。base 在逐字相同的配置下同一分歧可正常运行(eager 回退 rank 不发 host 集合、每层一次 router 调用与回放图配对)，故为本变更新引入的可用性回归。
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:151： dist.all_reduce(plan, op=dist.ReduceOp.MAX, group=dist.group.WORLD)
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:128： in COLLECTIVE_CHUNK_STRATEGIES
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:121： or (torch.cuda.is_available() and torch.cuda.is_current_stream_capturing())
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:152： token = SYNCHRONIZED_CHUNK_TOKENS.set(int(plan[0].item()))
rtp_llm/models_py/modules/dsv4/decode/forward.py:351： with synchronized_moe_chunk_plan(v4.layers, input_ids.numel(), input_ids.device):
rtp_llm/models_py/modules/dsv4/decode/forward.py:484： with synchronized_moe_chunk_plan(v4.layers, 0, param_dev):
rtp_llm/models_py/modules/dsv4/transformer.py:571： with synchronized_moe_chunk_plan(
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:108： in {"mega_moe", "mega_moe_se"}
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:109： and not getattr(getattr(layer, "ffn", layer), "is_decode_role", False)
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:235： or self.is_decode_role
rtp_llm/models_py/modules/dsv4/decode/forward.py:347： for layer_idx, layer in enumerate(v4.layers):
rtp_llm/models_py/modules/dsv4/decode/forward.py:470： return PyModelOutputs(
rtp_llm/cpp/cuda_graph/cuda_graph_decode.cc:6： replayGraph(bs);
rtp_llm/cpp/cuda_graph/cuda_graph_runner.cc:1963： & { graph_instances[key].graph.replay(); },
rtp_llm/cpp/cuda_graph/cuda_graph_runner.cc:2041： auto py_outputs_obj = py_forward_method(inputs, attn_pyobj);
rtp_llm/cpp/models/PyWrappedModel.cc:1105： const bool can_run_graph = !capture_hidden_states && enable_cuda_graph && graph_runner != nullptr
rtp_llm/cpp/models/PyWrappedModel.cc:1136： auto outputs = py_forward_method(py_model_inputs, held_attn_pyobj_);
rtp_llm/cpp/cuda_graph/cuda_graph_runner.cc:1295： RTP_LLM_LOG_WARNING("decode batch size %d exceeds max captured %d, fallback to normal run",
rtp_llm/test/smoke/suites_cuda13.bzl:65：--enable_cuda_graph 1 --decode_capture_config '1,2,4,8,16' --act_type BF16 --tp_size 1 --dp_size 2 --ep_size 2 --moe_strategy mega_moe_se
rtp_llm/test/smoke/suites_cuda13.bzl:65：--enable_cuda_graph 1 --decode_capture_config '1,2,4,8,16' --act_type BF16 --tp_size 1 --dp_size 2 --ep_size 2 --moe_strategy mega_moe_se
rtp_llm/config/engine_config.py:247： runtime_config.max_generate_batch_size = concurrency_config.concurrency_limit
rtp_llm/server/server_args/concurrent_group_args.py:22： default=32,
rtp_llm/cpp/engine_base/schedulers/FIFOScheduler.cc:145： if (running_streams_.size() + admitted_stream_count + 1 <= max_generate_batch_size_) {
rtp_llm/cpp/normal_engine/NormalEngine.cc:840： mayAddFakeStream(streams);
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_dp_router.py:12：In DP mode, different ranks may have different batch sizes (real requests vs
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/mega_moe.py:490： # every rank in buf.group MUST enter the kernel together. The kernel
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/mega_moe.py:509： # flow depending on a GPU-side scalar -> CUDA-graph-capture safe.
rtp_llm/models_py/distributed/collective_torch.py:233： infinite_timeout = timedelta(days=36500)
rtp_llm/models_py/modules/dsv4/transformer.py:615： # the missing rank's absence deadlocks the buffer with "CPU
rtp_llm/models_py/modules/dsv4/chunk_env.py:15： return os.environ.get("DSV4_MOE_CHUNK_PREFILL", "1") != "0"
rtp_llm/models_py/modules/dsv4/block.py:184： chunking_enabled=chunked_moe_enabled(),
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_fp4.py:55： strategy_name = "mega_moe_se"
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/mega_moe_se.py:60： def check_conditions(cls, checker: Any, config) -> None:
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/test/mega_moe_routed_only_execution_test.py:342： with torch.inference_mode(), torch.cuda.graph(graph):
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:112： is decoding while another rank is prefilling must execute the same number
rtp_llm/cpp/models/PyWrappedModel.h:407： graph_params.max_context_batch_size = params.concurrency_config.concurrency_limit;
建议：使计划集合与图模式对称或互斥：enable_cuda_graph 时 decode 侧跳过 forward 级计划归约(对齐 pure_dp 路由对图的拒绝方式)；或把计划改为固定形状设备端缓冲纳入捕获、由 C++ 在每次回放前代表回放 rank 执行同形归约；或将 tokens/chunk 协商并入路由自身已入图的集合以消除独立 WORLD 操作；或将图/eager 决策每步广播到全 WORLD 保证同一步全 rank 同臂。
评审版本：f8c79ba72870
| and config.enable_cuda_graph | ||
| and config.ep_size == 1 | ||
| and config.tp_size == 1 | ||
| and 9 <= get_sm()[0] < 12 |
There was a problem hiding this comment.
[P2] auto+CUDA Graph 兜底区间含 SM 10.x，使仓内未经图回放验证的 masked V2 成为默认选择（base 同配置为建模期 fail-fast）
SM 10.x（B200/GB200）单卡：tp_size=1、ep_size=1、dp_size=1，FP8_PER_BLOCK MoE 权重（如 Qwen3-30B-A3B FP8 blockwise 单卡部署），MOE_STRATEGY 保持默认 auto，use_all_gather 默认 True，用户显式 --enable_cuda_graph 1 并配置 decode_capture_config。GenericMoeLayer 建模时选中 masked V2，解码图捕获并回放其 execute_masked。base 同一配置在建模期即抛 “No suitable MOE strategy found”。
base 的该 check_conditions 只有 checker.check(config.moe_strategy == "fp8_per_block_no_dp_masked")（base fp8_per_block.py:66）；head 改为显式名或 auto_graph_fallback，并把架构区间写到 9<=get_sm()[0]<12（head fp8_per_block.py:68-78）。execute_masked、V2 的 check_conditions、hybrid 图守卫在 base/head 完全相同，差异只在选择层。
head 的 CudaFp8PerBlockNoDPMaskedStrategy.check_conditions 新增 auto_graph_fallback（auto + enable_cuda_graph + ep_size==1 + tp_size==1 + 9<=get_sm()[0]<12）。同场景下 DeepGemmHybridExecutor 的守卫 not enable_cuda_graph or sm==12（base/head 同为 204 行）淘汰 NoDP；DeepEP 路由要求 ep>1；Cutlass FP8 执行器只接受 per-tensor 量化；BatchedTriton 路由要求无量化；DeepGemmMaskedExecutorV2 只要求 sm>=9、无图与架构上限，因而成为唯一候选并被注册表选中（head 选择测试显式断言 SM=10 选中 masked）。引擎级 enable_cuda_graph 使 PyWrappedModel 创建解码 CudaGraphRunner（该分支无 MoE 架构门），捕获体在 graphCaptureBegin 之后执行 python forward，把 execute_masked 纳入图；SM 10.x 上该路径还走 UE8M0 packed-scale 专属分支，核集合与 SM90 已验证配置不同。
若 SM 10.x 上 masked V2（deep_gemm masked 分组 GEMM 与 SM100 UE8M0 packed-scale 分支）在捕获/回放下与 eager 不等价：轻则预热或捕获期异常导致启动失败，重则回放产出静默错误的解码结果并持续带病服务；base 对同配置为启动期 fail-fast，不会输出错误结果。若回放等价，则本改动只是把既有显式路径变为默认可达（预期收益）。
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_per_block.py:73： and 9 <= get_sm()[0] < 12
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_per_block.py:69： config.moe_strategy == "auto"
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_hybrid_executor.py:204： checker.check(not config.enable_cuda_graph or get_sm()[0] == 12)
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_hybrid_executor.py:204： checker.check(not config.enable_cuda_graph or get_sm()[0] == 12)
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_masked_executor_v2.py:41： checker.check(get_sm()[0] >= 9)
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_per_block.py:66： checker.check(config.moe_strategy == "fp8_per_block_no_dp_masked")
rtp_llm/models_py/modules/factory/fused_moe/tests/test_cuda_strategies.py:386： (10, CudaFp8PerBlockNoDPMaskedStrategy),
rtp_llm/cpp/config/ConfigModules.h:342： std::string moe_strategy = "auto";
rtp_llm/cpp/config/ConfigModules.h:302： bool enable_cuda_graph = false;
rtp_llm/models_py/modules/factory/fused_moe/strategy_registry.py:174： "No suitable MOE strategy found for configuration. "
rtp_llm/models_py/modules/factory/fused_moe/utils/config_resolver.py:76： return config.ep_size > 1
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/test/deepgemm_masked_executor_v2_cuda_graph_test.py:40： "masked DeepGEMM CUDA graph test requires SM90",
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/test/BUILD:40： tags = ["H20_CU13"],
rtp_llm/cpp/models/GenerationPrefillCudaGraphEligibility.h:92：// architecture passes the masked-MoE graph replay test.
rtp_llm/cpp/models/GenerationPrefillCudaGraphEligibility.h:94： return cuda_compute_capability_major == 9;
rtp_llm/cpp/models/test/ModelDataTest.cc:580： EXPECT_FALSE(supportsGenerationPrefillCudaGraphMaskedMoeBackend(10));
rtp_llm/cpp/models/PyWrappedModel.h:472： std::make_unique(graph_params, py_instance, forward_method, params.metrics_reporter);
rtp_llm/cpp/cuda_graph/cuda_graph_runner.cc:2041： auto py_outputs_obj = py_forward_method_(inputs, attn_pyobj);
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_hybrid_executor.py:574： sm_major == 10
rtp_llm/test/smoke/suites_sm100.bzl:101： # REMOVED: moe_masked_cudagraph_tp2_sm100 (Masked MoE + FP8_PER_BLOCK + CudaGraph + TP2)
rtp_llm/models_py/model_desc/generic_moe.py:88： enable_cuda_graph=enable_cuda_graph,
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_tp_router.py:66： checker.check(resolver.use_all_gather(config))
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/deepep_normal_router.py:48： checker.check(resolver.is_ep_enabled(config))
rtp_llm/models_py/modules/factory/fused_moe/impl/common/router/batched_data_router.py:50： checker.check(not resolver.has_quantization(config))
rtp_llm/server/server_args/moe_group_args.py:187： default=True,
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/cutlass_moe.py:54： quant_method in ["FP8_PER_TENSOR_COMPRESSED", "FP8_DYNAMIC_PER_TENSOR"]
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/test/deepgemm_masked_executor_sm100_test.py:26： self._test_deepgemm_masked_executor(True)
待确认：SM 10.x（B200/GB200）真机上 deep_gemm masked 分组 GEMM 与 UE8M0 packed 核在 CUDA Graph 捕获/回放下的行为：3rdparty/deep_gemm 在快照内只有两个补丁与 BUILD，依赖源码缺失；仓内也没有 SM 10.x 的 masked V2 图测试或端到端用例。该未知直接决定本缺陷是否成立。
建议：把兜底架构区间收紧为 get_sm()[0] == 9（与 C++ 掩码 MoE 图后端的按架构门一致），SM 10.x 暂仍要求显式 --moe_strategy fp8_per_block_no_dp_masked；或先为 SM100 复制 deepgemm_masked_executor_v2_cuda_graph_test.py 的捕获+更新路由回放+eager 等价校验（去掉 SM90 断言、绑定 SM100_ARM_CU13），并在 SM100 冒烟矩阵恢复 masked+图用例，通过后再放开 10<=sm<12。
评审版本：f8c79ba72870
| # ``model.layers.*``. The original RTP conversion emits split FP4 | ||
| # experts under ``layers.*``. Keep both layouts available and let the | ||
| # runtime quantization setting decide which kernels consume them. | ||
| self._raw_hf_stacked = any( |
There was a problem hiding this comment.
[P1] V4 raw HF 布局标志 _raw_hf_stacked 被 finetune 元数据的第二次 _process_meta 无条件覆盖,raw 检查点叠加 ptuning 时全部描述符键翻转导致加载确定性失败
部署 V4 raw HF 检查点(safetensors 键含 model.layers..mlp.experts.gate_up_proj,首次 meta 调用置 True),同时设置通用参数 --ptuning_path 指向含 pytorch_model.bin 的目录:该文件进入 finetune_file_list,第二次 _process_meta 的键集不含标记(前缀微调文件不含 V4 专家键;即使元数据为空,守卫按文件计数仍会调用),标志翻回 False。
base 的 _key 无实例状态、恒返回 converted 前缀(base deepseek_v4.py:181),base _process_meta 无布局标志;head 引入 _raw_hf_stacked 且仅在被多次调用的 _process_meta 中按当次键集赋值,_key/routed expert/全局键全面依赖该状态。
_process_meta 用 self._raw_hf_stacked = any(...) 对当次调用的键集直接赋值且无构造期/类级默认;create_model_weight_info 先以 pretrain 键集调用一次(置 True),再以仅含 ptuning 文件的 finetune 键集调用一次,process_meta_from_ckpt 的早退守卫只看文件数不看键集,故第二次必然执行并把标志覆盖为 False;随后 get_weight_info 生成的 _key、routed expert 与全局键全部切换到 converted 布局键名。
所有描述符改用 converted 布局键(layers.、FP4 拆分专家、embed.weight 等),raw 检查点均不存在;load_tensor 对缺名返回空列表,identity 抛 'ts is empty',AtomicWeight 记录日志后重抛,模型加载确定性失败。属响亮阻塞失败,无静默错权;新功能宣称的 raw 布局支持在与框架通用微调机制组合时失效。
rtp_llm/models/deepseek_v4.py:159：self._raw_hf_stacked = any(
rtp_llm/models/deepseek_v4.py:160：".mlp.experts.gate_up_proj" in key for key in weight_keys
rtp_llm/models/deepseek_v4.py:234：if self._raw_hf_stacked:
rtp_llm/models/deepseek_v4.py:594：if self._raw_hf_stacked:
rtp_llm/models/deepseek_v4.py:615：if self._raw_hf_stacked
rtp_llm/model_loader/model_weight_info.py:627：self.process_meta_from_ckpt(database.pretrain_file_list)
rtp_llm/model_loader/model_weight_info.py:628：self.process_meta_from_ckpt(database.finetune_file_list)
rtp_llm/model_loader/model_weight_info.py:629：weight_info = self.get_weight_info()
rtp_llm/model_loader/model_weight_info.py:706：if len(ckpt_metas) == 0:
rtp_llm/utils/database.py:214：for f in Path(ptuning_path).glob("pytorch_model.bin"):
rtp_llm/utils/database.py:219：self.finetune_file_list.append(ckpt)
rtp_llm/utils/database.py:111：self._tensor_index[tname] = ckpt_file
rtp_llm/utils/database.py:154：Finetune files, such as ptuning weights, are intentionally left untouched
rtp_llm/models/base_model.py:492：self.model_config.ptuning_path,
rtp_llm/server/server_args/model_group_args.py:57："--ptuning_path",
rtp_llm/utils/database.py:244：return []
rtp_llm/utils/model_weight.py:167：raise Exception("ts is empty")
rtp_llm/model_loader/weight_module.py:380：raise e
rtp_llm/models/deepseek_v4.py:181：return f"{self._LAYER_PREFIX}.{suffix}"
rtp_llm/models/qwen3_next/qwen3_next_weight.py:560：self._has_stacked_ckpt = False
rtp_llm/models/qwen3_next/qwen3_next_weight.py:575：self._has_stacked_ckpt = True
rtp_llm/model_loader/loader.py:59：self._weights_info.create_model_weight_info(database)
rtp_llm/models/deepseek_v4.py:155：# Fine-tuned HF exports retain stacked BF16 experts under
rtp_llm/models/test/moe_cuda13_config_propagation_test.py:585：builder._process_meta(
rtp_llm/models/deepseek_v41.py:26：self._has_vision_router_bias = any(
建议：构造期初始化 _raw_hf_stacked=False,并在 _process_meta 中粘滞置位(self._raw_hf_stacked = self._raw_hf_stacked or any(...)),与 qwen3_next_weight.py:560/575 的既有惯例一致;或合并 pretrain+finetune 键集后一次性判定,同时为未经 meta 的直接构造提供默认值。
评审版本：f8c79ba72870
| @@ -1015,12 +1032,17 @@ def _load_raw_tensor( | |||
| ) | |||
|
|
|||
| is_dense_weight = self.kernel.name not in (W.moe_w1, W.moe_w2) | |||
There was a problem hiding this comment.
[P2] 非 e8m0(SM90/91)在线量化的 legacy-reshape 豁免仅覆盖 wo_a:其余 V4 2D 权重以 [K,N]+fp32 入库违反消费端 [N,K] 契约,非方形共享几何下装载期必崩;tp>1 错轴切片被 base/head 同文的既有 workspace 失配掩蔽
SM90/91(capability 主号 9,is_deep_gemm_e8m0_used()==False;SM89 因 DeepGemmHybridExecutor 的 get_sm()[0]>=9 先行排除)+ 原生 HF BF16 V4 导出(含 .mlp.experts.gate_up_proj 键)+ --quantization FP8_PER_BLOCK:(1)共享几何非方形(dim≠2moe_inter_dimn_shared_experts,配置解析只校验 n_shared 为非负整数,head 自测夹具 hidden 512/moe_inter 256/n_shared=2 即非方形):任意 tp(含 tp=1)在 Block→ChunkedFp8Fp4MoeLayer→Fp8Fp4MoeLayer→W13SharedExpert 构造期 RuntimeError,先于任何 GEMM(mega_moe_se 的 includes_shared_expert=True 分支在非 e8m0 被 capability 门禁排除,必走 W13SharedExpert);(2)注意力 tp>1:错轴分片在装载期即物化,但该配置在 base/head 均先因既有失配崩溃——workspace q_dim 以全局头数绑定(head:660 与 base:655 逐字同文,n_heads=attn_config.head_num),prefill_q 返回全局宽度视图,forward 以本地头数 view(head:5361/base:5330),任意设备首次 prefill 即 RuntimeError;(3)tp=1+方形几何(快照已知 V4 几何 dim=4096=220481):head 全链自洽可用,相对 base 为修复。
base 无 raw-HF 键映射(_key 恒为 layers.{i} 模板),原生 BF16 导出+在线量化在 base 装载期即失败;base 对在线量化一律 .T(base:1058)再 reshape,标签 [N,K] 而数据为 W^T 扁平重排,非 e8m0 下 linear 读得 self.K=N、self.N=K,首次前向 _validate_input 必崩(方形则静默错误)。head 新增 _raw_hf_stacked/_RAW_SUFFIXES 使该组合首次可达,layer.py 新增报错文本明确指引『Raw BF16 checkpoints require --quantization FP8_PER_BLOCK』;head 修复了 e8m0 设备(direct_ue8m0→[N,K]+int32)、wo_a(raw_ue8m0 旁路)与 tp=1 稠密链路(精确往返),但未把非 e8m0 豁免扩展到同判据(is_v4_weight)的其余 V4 权重。
head 的 _load_raw_tensor 对 is_v4_weight 置 preserve_checkpoint_layout 跳过 kernel/scale 的 .T,非 e8m0(is_deep_gemm_e8m0_used 仅认 capability 主号 10/12)分支经 per_block_cast_to_fp8 得到内存 [N,K] 的 fp8 权重与 [N/128,K/128] fp32 标度;但 LoadQuantPerBlockFp8Weight._postprocess 的豁免条件只匹配 W.v4_attn_wo_a_w+float8_e8m0fnu,其余 V4 2D 权重落入 PerBlockFp8Weight._postprocess 的 kernel/scale 各一次 reshape(shape[-1],-1),声明标签翻转为 [K,N]+[K/128,N/128](内存逻辑序仍为 [N,K])。dsv4 消费端(base/head 同构)按 [N,K] 契约工作:_fp8_w_s 行切 dim0=N、非 int32 标度按 //128 块切、wo_b 列切 dim1=K;W13SharedExpert 构造期强制 (2inter_dim,dim);装载侧对 V4 权重登记 sp_id 把分片职责交给模块。tp=1 时 CudaFp8DeepGEMMLinear 非 e8m0 分支按 (K,N) 读标签再 reshape(N,K) 精确复原、fp32 标度校验通过,链路自洽;共享 w13 在非方形几何下真实 [2inter,dim] 被标为 [dim,2*inter],构造期立即 RuntimeError;tp>1 时行/列切片作用于翻转后的错轴(rank0 clamp 为整块、rank>0 空分片,构造期 fp32 校验对两者恒过,含 K=0),但该配置在 base/head 均先崩溃于同文的 workspace 失配。
若 SM90/91 属 DSV4 受支持部署:head 自身文本指引的 raw-BF16+FP8_PER_BLOCK 配置在非方形共享几何下确定性地无法完成模型构造(装载期 RuntimeError『shared w13 weight shape mismatch』),服务无法启动;同一检查点在 e8m0 主线可正常装载,失败特定于本 PR 未豁免的非 e8m0 分支,且 base 对该配置同样不可用(base 首次前向崩溃),故属新启用路径交付不完整而非既有可用配置的回归。tp>1 腿为潜在契约不一致:两版在同一 view 处先崩,一旦 workspace 失配修复即在非 e8m0 上暴露为分片错位(rank0 计算整块未分片输出=错误结果,rank>0 空分片)。均为响亮失败;head 下无新增静默错误输出(方形恒等重标数据正确,base 的方形静默转置反被 head 修复)。
rtp_llm/model_loader/per_block_fp8_quant_weight.py:1034： is_dense_weight = self.kernel.name not in (W.moe_w1, W.moe_w2)
preserve_checkpoint_layout = is_v4_weight(self.kernel)
rtp_llm/model_loader/per_block_fp8_quant_weight.py:1096： elif (
quant_kernel.dim() == 2
and not direct_ue8m0
and not preserve_checkpoint_layout
):
quant_kernel = quant_kernel.T
rtp_llm/model_loader/per_block_fp8_quant_weight.py:866： if not is_deep_gemm_e8m0_used():
kernel_weight = (
kernel_weight.reshape(kernel_weight.shape[-1], -1)
if kernel_weight.dim() == 2
else kernel_weight
)
rtp_llm/model_loader/per_block_fp8_quant_weight.py:875： if not is_deep_gemm_e8m0_used():
scale_weight = (
scale_weight.reshape(scale_weight.shape[-1], -1)
if scale_weight.dim() == 2
else scale_weight
)
rtp_llm/model_loader/per_block_fp8_quant_weight.py:957： if (
self.kernel.name == W.v4_attn_wo_a_w
and self.scale is not None
and isinstance(tensor, dict)
and tensor[self.scale.name].dtype == torch.float8_e8m0fnu
):
# The V4 wo_a einsum consumes checkpoint-layout FP8 and raw UE8M0
# block scales; the legacy FP8 linear reshape would corrupt both.
return CompositeWeight._postprocess(self, tensor, device, load_config)
rtp_llm/model_loader/per_block_fp8_quant_weight.py:1038： raw_wo_a_ue8m0 = (
self.scale is not None
and self.kernel.name == W.v4_attn_wo_a_w
and not e8m0_used
)
rtp_llm/model_loader/per_block_fp8_quant_weight.py:1287： return CompositeWeight._postprocess(self, tensor, device, load_config)
rtp_llm/model_loader/per_block_fp8_quant_weight.py:214： # ---- DSv4 (TP=1 placeholders) ----
W.v4_attn_wq_a_w: sp_id,
W.v4_attn_wq_a_s: sp_id,
W.v4_attn_wq_b_w: sp_id,
W.v4_attn_wq_b_s: sp_id,
rtp_llm/model_loader/per_block_fp8_quant_weight.py:1157：PerBlockFp8Weight.w8a8_weight_list.update(_V4_FP8_WEIGHT_LIST)
rtp_llm/models_py/kernels/cuda/deepgemm_wrapper.py:92：def is_deep_gemm_e8m0_used() -> bool:
# Blackwell SM100 and SM120 DeepGEMM kernels consume packed UE8M0 scales.
# SM120 support requires the vLLM-pinned DeepGEMM build (a6b593d or newer).
return torch.cuda.get_device_capability()[0] in (10, 12)
rtp_llm/model_loader/per_block_fp8_quant_weight.py:1056： if self.kernel.name == W.moe_w1 or self.kernel.name == W.moe_w2:
pass
elif quant_kernel.dim() == 2 and not direct_ue8m0:
quant_kernel = quant_kernel.T
rtp_llm/models/deepseek_v4.py:159： self._raw_hf_stacked = any(
".mlp.experts.gate_up_proj" in key for key in weight_keys
)
rtp_llm/models/deepseek_v4.py:233： def _key(self, suffix: str) -> str:
if self._raw_hf_stacked:
rtp_llm/models/deepseek_v4.py:180： def _key(self, suffix: str) -> str:
return f"{self._LAYER_PREFIX}.{suffix}"
rtp_llm/models/deepseek_v4.py:426： AttnAtomicWeight(
W.v4_shared_w13_w,
[
CkptWeightInfo(
self._key("ffn.shared_experts.w1.weight"),
identity,
),
CkptWeightInfo(
self._key("ffn.shared_experts.w3.weight"),
identity,
),
],
concat_0,
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:236： raise RuntimeError(
"DeepSeek-V4 routed experts must be FP8-per-block or packed FP4; "
f"got {routed_dtype}. Raw BF16 checkpoints require "
"--quantization FP8_PER_BLOCK."
)
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:292： if self.fused_moe.includes_shared_expert or n_shared_experts == 0:
self.shared_experts = None
self._shared_executor = None
else:
shared_inter_dim = moe_inter_dim * n_shared_experts
self.shared_experts = W13SharedExpert(
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/shared_expert.py:139： if w13_w.shape[0] != 2 * inter_dim or w13_w.shape[1] != dim:
raise RuntimeError(
"shared w13 weight shape mismatch: "
f"got {tuple(w13_w.shape)}, expected {(2 * inter_dim, dim)}"
)
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/weight_adapter.py:56： shared_gate_up_name = weight_names.get("shared_gate_up")
if n_shared_experts > 0 and shared_gate_up_name and shared_gate_up_name in weights:
weights[W.ffn_w13] = weights.pop(shared_gate_up_name)
rtp_llm/models_py/modules/dsv4/block.py:157： adapt_split_moe_weights(
layer_weights, moe_inter_dim, n_shared_experts, _MOE_WEIGHT_NAMES
)
rtp_llm/models_py/modules/dsv4/fp8/attention.py:963： scale_is_packed_int32 = s.dtype == torch.int32
if row_slice is not None:
w = w[row_slice]
if scale_is_packed_int32:
s = s[row_slice]
else:
s = s[row_slice.start // 128 : row_slice.stop // 128]
rtp_llm/models_py/modules/dsv4/fp8/attention.py:988： self.wq_b = _fp8_w_s(
W.v4_attn_wq_b_w,
W.v4_attn_wq_b_s,
row_slice=wq_b_row_slice if tp_size > 1 else None,
)
rtp_llm/models_py/modules/factory/linear/impl/cuda/fp8_deepgemm_linear.py:87： else:
# Reshape weight and weight scale
self.K, self.N = self.weight.shape
self.scale_K, self.scale_N = self.weight_scales.shape
self.weight = self.weight.reshape(self.N, self.K)
self.weight_scales = self.weight_scales.reshape(self.scale_N, self.scale_K)
rtp_llm/models_py/modules/factory/linear/impl/cuda/fp8_deepgemm_linear.py:94： if self.weight_scales.dtype is torch.float32 and (
(self.N + 127) // 128 != self.scale_N
or (self.K + 127) // 128 != self.scale_K
):
rtp_llm/models_py/modules/factory/linear/impl/cuda/fp8_deepgemm_linear.py:186： if K != self.K:
error_msg = f"Input tensor inner dimension expected to be {self.K}, got {K}"
logger.error(error_msg)
raise ValueError(error_msg)
rtp_llm/models_py/model_desc/deepseek_v4_model.py:660： q_dim = int(self._v4_args.n_heads) * int(self._v4_args.head_dim)
rtp_llm/models_py/model_desc/deepseek_v4_model.py:655： q_dim = int(self._v4_args.n_heads) * int(self._v4_args.head_dim)
rtp_llm/models_py/modules/dsv4/fp8/attention.py:5357： q_out = common.workspace.prefill_q(seqlen)
q_local_flat = self._lin(
self.wq_b,
qr_3d,
out=q_out.view(seqlen, self.n_heads * self.head_dim),
)
rtp_llm/models_py/modules/dsv4/fp8/attention.py:5326： q_out = common.workspace.prefill_q(seqlen)
q_local_flat = self._lin(
self.wq_b,
qr_3d,
out=q_out.view(seqlen, self.n_heads * self.head_dim),
)
rtp_llm/models_py/model_desc/deepseek_v41_base_model.py:544： q_dim = int(v4.layers[0].attn.n_heads) * int(self._v4_args.head_dim)
rtp_llm/model_loader/test/test_stacked_moe_weight.py:183： with patch(
"rtp_llm.models_py.kernels.cuda.deepgemm_wrapper.is_deep_gemm_e8m0_used",
return_value=False,
):
result = wrapped.load(source, 0, "cpu", load_config)
rtp_llm/model_loader/test/test_stacked_moe_weight.py:191： self.assertEqual(tuple(weight.shape), (256, 512))
self.assertEqual(weight.dtype, torch.float8_e4m3fn)
self.assertEqual(tuple(scale.shape), (2, 4))
self.assertEqual(scale.dtype, torch.float8_e8m0fnu)
rtp_llm/model_loader/test/test_stacked_moe_weight.py:111： with patch(
"rtp_llm.models_py.kernels.cuda.deepgemm_wrapper.is_deep_gemm_e8m0_used",
return_value=False,
), patch(
"rtp_llm.model_loader.per_block_fp8_quant_weight.per_block_cast_to_fp8",
return_value=(quantized, scale),
):
result = wrapped._load_raw_tensor(source, 0, "cpu", load_config)
rtp_llm/model_loader/test/BUILD:11： exec_properties = {"gpu": "L20_CU13"},
tags = ["L20_CU13"],
rtp_llm/model_loader/test/BUILD:71： # Metadata-only test; L20_CU13 supplies the runtime without an H20 reservation.
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_hybrid_executor.py:200： checker.check(quant_method == "FP8_PER_BLOCK")
checker.check(resolver.is_bf16(config))
checker.check(has_deep_gemm())
checker.check(get_sm()[0] >= 9)
checker.check(not config.enable_cuda_graph or get_sm()[0] == 12)
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_per_block.py:68： auto_graph_fallback = (
config.moe_strategy == "auto"
and config.enable_cuda_graph
and config.ep_size == 1
and config.tp_size == 1
and 9 <= get_sm()[0] < 12
)
rtp_llm/models_py/modules/dsv4/test/defs.bzl:9： cuda13_sm100_py_test(
rtp_llm/models_py/modules/dsv4/transformer.py:79： moe_inter_dim: int = 2048
n_routed_experts: int = 256
n_physical_experts: Optional[int] = None
n_shared_experts: int = 1
rtp_llm/models_py/modules/dsv4/transformer.py:52： dim: int = 4096
rtp_llm/models/test/moe_cuda13_config_propagation_test.py:213： "moe_intermediate_size": 256,
"n_shared_experts": n_shared_experts,
rtp_llm/models/test/moe_cuda13_config_propagation_test.py:197： "hidden_size": 512,
rtp_llm/models/test/moe_cuda13_config_propagation_test.py:494： model_config = self._load_minimal_model_config(2)
rtp_llm/models/deepseek_v4.py:126： if isinstance(raw_value, bool) or not isinstance(raw_value, int) or raw_value < 0:
raise ValueError(
"DeepSeek-V4 config.json field 'n_shared_experts' must be a "
f"non-negative integer, got {raw_value!r}"
)
rtp_llm/models_py/modules/dsv4/fp8/decode/mega_support.py:159： if capability not in SUPPORTED_CAPABILITIES:
return (
f"sm_100a or sm_103a is required, got sm{capability[0]}{capability[1]}",
None,
)
rtp_llm/utils/model_weight.py:1756：def is_v4_weight(src_weight_info) -> bool:
"""Return True iff src_weight_info.name belongs to the DSv4 namespace
(v4.…). Used by V4-specific quant subclass support() checks so they
do not accidentally match DSv2/V3 weights with similar shapes."""
name = getattr(src_weight_info, "name", None)
return isinstance(name, str) and name.startswith("v4.")
rtp_llm/models_py/modules/dsv4/fp8/attention.py:399： assert scale_raw.dtype == torch.float8_e8m0fnu, (
f"unexpected raw wo_a scale dtype {scale_raw.dtype}"
)
rtp_llm/models_py/modules/dsv4/prefill_workspace.py:132： return (
self._union[: self._q_bytes]
.view(torch.bfloat16)
.view(self._q_rows, self._q_dim)[:num_tokens]
)
rtp_llm/models_py/model_desc/deepseek_v4_model.py:289： n_heads=attn_config.head_num,
rtp_llm/models_py/modules/dsv4/fp8/attention.py:933： if row_slice is not None:
w = w[row_slice]
if scale_is_packed_int32:
s = s[row_slice]
else:
s = s[row_slice.start // 128 : row_slice.stop // 128]
待确认：SM90/91(Hopper)是否为 DSV4『原生 HF BF16 + --quantization FP8_PER_BLOCK 在线量化』的受支持部署目标:快照内既无支持矩阵声明也无装载/启动期拒绝,正向证据为 head 自建的 raw_wo_a_ue8m0 旁路+corrupt 注释+非 e8m0 全量 load() 测试+L20/H20 池注释+sm>=9 执行器与 9<=sm<12 策略分支+无平台限定的指引文本,反向证据为 V4 smoke 仅 sm100/sm103、dsv4 测试全绑 SM100 池、deep_gemm 在 SM90 的运行期内核行为不在快照内;该前提成立与否直接决定装载期崩溃是否发生在受支持路径上
建议：把 LoadQuantPerBlockFp8Weight._postprocess 的豁免判据从 wo_a+e8m0fnu 扩展为 is_v4_weight(self.kernel)(与 preserve_checkpoint_layout 同判据),非 e8m0 统一交付 [N,K] 并让标度与消费端适配(wo_a 式 e8m0fnu,或让 _fp8_w_s/W13SharedExpert/create_fp8_linear 显式识别 [K,N]+fp32 转置标签);或在不支持的平台/几何组合上于装载期显式 fail-fast。同时补『非 e8m0 在线量化 + 非方形共享几何 + Block 构造』与『非 e8m0 + tp>1』的全管线集成测试(现有 wq_a 非 e8m0 测试止步 _load_raw_tensor,未覆盖 _postprocess 翻转);并修复或显式拒绝 deepseek_v4_model.py:660 的全局 q_dim 与本地头 view 失配(可参照 deepseek_v41_base_model.py:544),避免其继续掩蔽 V4 的 TP 分片缺陷。
评审版本：f8c79ba72870
| plan = torch.tensor( | ||
| [int(tokens), local_chunk_count], dtype=torch.int64, device=device | ||
| ) | ||
| dist.all_reduce(plan, op=dist.ReduceOp.MAX, group=dist.group.WORLD) |
There was a problem hiding this comment.
[P1] head 把计划 all_reduce 引入 eager 解码侧而 CUDA 图回放不参与:mega_moe(_se)+enable_cuda_graph+dp>1 的 DECODE 配置下,rank 本地图/eager 回退分歧使 eager rank 的 NCCL 计划 all_reduce 无对端、回放 rank 的 mega NVLink 屏障超时 trap,解码实例挂死并崩溃(base 对同一分歧对称无害,属本次引入的回归)
仓库固化的受支持配置(head 与 base 相同的 ARM smoke decode 角色,suites_cuda13.bzl:65):--enable_cuda_graph 1 --decode_capture_config '1,2,4,8,16' --tp_size 1 --dp_size 2 --ep_size 2 --moe_strategy mega_moe_se --role_type DECODE --sp_type dspark。某解码步 rank X 本地 decode bs>16(cuda_graph_runner.cc:1295 显式记录该回退)或 canReplaySelectedGraph 的数据相关校验(如 cuda_graph_runner.cc:1381 的 combo_position_ids 校验)仅在 rank X 失败 → rank X 走 PyWrappedModel.cc:1136 的 eager Python 解码臂并进入计划 all_reduce;rank Y 命中捕获桶 → graph_runner->forward 纯本地图回放,无对应 NCCL 调用。两 rank 同处一个 WORLD 组,DP 批次相互独立(pure_dp_router.py:12-13),空 rank 亦由假流进入同一步(NormalEngine.cc:954)。
base 的计划只包裹 prefill 臂(prefill/forward.py:622、transformer.py:683),needs_sync 显式排除 decode-role 层(base chunked_layer.py:109),base 解码臂与 forward_decode 均无计划包裹(base transformer.py:563 为裸层循环、base decode/forward.py:469 空批直接短路返回),base 解码不分块(base chunked_layer.py:216 分支返回 False),故 eager 与回放都是每层 1 次 mega 调用且全程无 NCCL 调用,rank 间图/eager 分歧在 base 对称无害。head 去掉角色排除、新增 SYNCHRONIZED_CHUNK_COUNT 并把计划 all_reduce 加到所有 eager 解码侧前向(decode/forward.py:351/:484、transformer.py:571),使同一分歧变为单边 WORLD NCCL 调用而死锁。混合 prefill/decode(PDFUSION)场景 base 已在计划 all_reduce 层面失配,该暴露非本次引入,不在本发现范围。
head 的解码侧 eager 前向(decode 批经 decode/forward.py:351,空 DP 批经 :484,forward_decode 经 transformer.py:571)在集体分块策略下进入 synchronized_moe_chunk_plan:needs_sync 只检查 strategy_name 是否属于 COLLECTIVE_CHUNK_STRATEGIES 与 chunking_enabled(head 已无 decode-role 排除),mega_moe_se 命中且 chunking_enabled 由 chunked_moe_enabled() 默认为真(decode 环境未设 DSV4_MOE_CHUNK_PREFILL/DSV4_CHUNK_TOKENS),于是每次 eager 解码前向在层循环之前执行一次 dist.all_reduce(plan, group=WORLD) 并在 plan[0].item() 上 CPU 同步阻塞。CUDA 图捕获时 capturing 守卫令计划直接 yield(捕获图内不含该 all_reduce),回放路径只调用 graph_runner->forward,内部 replayGraph 即 graph.replay() 的纯本地图重放、不执行任何 Python,故回放 rank 永不发起对应 NCCL 调用。graph/eager 选择是 rank 本地 canRun 决策,回退是显式支持并带日志/计数的运行路径(decode bs 超过捕获上限、combo_position_ids 校验失败、tagged-kv 不匹配),且 DP 各 rank 批次与元数据相互独立(空 rank 由最小假流参与同一步)。一旦同一解码步 rank X 回退 eager、rank Y 图回放:rank X 阻塞在无对端的 all_reduce,永不进入 mega 内核;rank Y 回放的 mega 内核是对等对称 NVLink 集合,缺 rank X 参与即屏障超时触发设备断言 trap。
跨 rank 双向死锁:rank X 的计划 all_reduce 无对端,阻塞至 NCCL 看门狗超时;rank Y 回放的 mega 内核 NVLink 屏障等不到 rank X,超时后 DG_DEVICE_ASSERT 触发 kernel 侧 trap,表现为 SIGTRAP/CUDA_ERROR_LAUNCH_FAILED。decode 角色环境未设 DG_MEGA_MOE_NVLINK_BARRIER_TIMEOUT_SECS(该变量只出现在 prefill 环境,suites_cuda13.bzl:93),按 mega_moe.py:502 注释约 30 秒崩溃。整个解码实例不可用,属可用性严重故障;base 上同一回退场景可正常运行。
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:151：dist.all_reduce(plan, op=dist.ReduceOp.MAX, group=dist.group.WORLD)
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:152：token = SYNCHRONIZED_CHUNK_TOKENS.set(int(plan[0].item()))
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:121：or (torch.cuda.is_available() and torch.cuda.is_current_stream_capturing())
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:128：in COLLECTIVE_CHUNK_STRATEGIES
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:112：is decoding while another rank is prefilling must execute the same number
rtp_llm/models_py/modules/dsv4/decode/forward.py:351：with synchronized_moe_chunk_plan(v4.layers, input_ids.numel(), input_ids.device):
rtp_llm/models_py/modules/dsv4/decode/forward.py:484：with synchronized_moe_chunk_plan(v4.layers, 0, param_dev):
rtp_llm/models_py/modules/dsv4/decode/forward.py:476：# Empty DP ranks still have to enter each routed-MoE collective
rtp_llm/models_py/modules/dsv4/transformer.py:571：with synchronized_moe_chunk_plan(
rtp_llm/models_py/modules/dsv4/transformer.py:615：the missing rank's absence deadlocks the buffer with "CPU
rtp_llm/models_py/modules/dsv4/block.py:184：chunking_enabled=chunked_moe_enabled(),
rtp_llm/models_py/modules/dsv4/chunk_env.py:15：return os.environ.get("DSV4_MOE_CHUNK_PREFILL", "1") != "0"
rtp_llm/cpp/models/PyWrappedModel.cc:1105：const bool can_run_graph = !capture_hidden_states && enable_cuda_graph && graph_runner != nullptr
rtp_llm/cpp/models/PyWrappedModel.cc:1123：py_model_outputs = graph_runner->forward(py_model_inputs, graph_state);
rtp_llm/cpp/models/PyWrappedModel.cc:1136：auto outputs = py_forward_method(py_model_inputs, held_attn_pyobj);
rtp_llm/cpp/cuda_graph/cuda_graph_runner.cc:1295：RTP_LLM_LOG_WARNING("decode batch size %d exceeds max captured %d, fallback to normal run",
rtp_llm/cpp/cuda_graph/cuda_graph_runner.cc:1381：if (!validateComboPositionIds(inputs, state, captured_position_ids, copy_numel)) {
rtp_llm/cpp/cuda_graph/cuda_graph_runner.cc:1963：& { graph_instances_[key].graph_.replay(); },
rtp_llm/cpp/cuda_graph/cuda_graph_decode.cc:33：capture_bs.push_back(max_generate_batch_size);
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/mega_moe.py:490：every rank in buf.group MUST enter the kernel together. The kernel
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/mega_moe.py:501：DG_DEVICE_ASSERT(false and "NVLink barrier timeout")
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_fp4.py:29：checker.check(config.moe_strategy in ("auto", cls.strategy_name))
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/fp8_fp4_router.py:65：checker.check(config.ep_size > 1)
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_dp_router.py:85：checker.check(not config.enable_cuda_graph)
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_dp_router.py:12：In DP mode, different ranks may have different batch sizes (real requests vs
rtp_llm/cpp/normal_engine/NormalEngine.cc:954：streams.emplace_back(MtpExecutor::createMinFakeDecodeStream(
rtp_llm/test/smoke/suites_cuda13.bzl:65：--enable_cuda_graph 1 --decode_capture_config '1,2,4,8,16' --act_type BF16 --tp_size 1 --dp_size 2 --ep_size 2 --moe_strategy mega_moe_se
rtp_llm/test/smoke/suites_cuda13.bzl:65：--enable_cuda_graph 1 --decode_capture_config '1,2,4,8,16' --act_type BF16 --tp_size 1 --dp_size 2 --ep_size 2 --moe_strategy mega_moe_se
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:109：and not getattr(getattr(layer, "ffn", layer), "_is_decode_role", False)
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:216：if is_decode or capturing:
rtp_llm/models_py/modules/dsv4/transformer.py:563：for layer_idx, layer in enumerate(self.layers):
rtp_llm/models_py/modules/dsv4/decode/forward.py:469：# Empty batch (B == 0) — short-circuit with zero-row hidden.
rtp_llm/models_py/model_desc/deepseek_v4_dspark_model.py:747：for layer_idx, layer in enumerate(self.v4.layers):
建议：使图与计划的参与面一致,任选其一:(1)在 canRun 决策后沿 DP 维度协商一致性(每步对 replay 标志做一次跨 rank 归约,任一 rank 回退则全体 eager);(2)仿照 pure_dp_router.py:85 的选择期检查,在集体分块策略(mega 系/fp8_per_block_ep_normal)启用且 dp>1 时拒绝 enable_cuda_graph 或强制全 eager;(3)以支持图捕获的 NCCL 通信域把计划 all_reduce 按固定次数纳入捕获,使回放 rank 同样携带对端调用。另可让 DSpARK 草稿循环复用 forward_layers 的计划包裹,消除契约声明与实现的偏差(当前草稿两侧均不进计划,暂无独立失配)。
评审版本：f8c79ba72870
| ) | ||
| args.parallelism_config = parallelism_config | ||
| args.moe_config = moe_config | ||
| args.enable_cuda_graph = bool( |
There was a problem hiding this comment.
[P2] auto 模式下单卡开启 CUDA graph 后，DSV4 同一 ffn 实例的 prefill 大 token 批恒走 DeepGemmMaskedExecutorV2 的 E×align(T,128) 无上限 masked 缓冲，长 prompt 可触发 CUDA OOM
单卡（ep==1 且 tp==1）FP8-per-block DSV4（快照注释记录的 v4-flash 1×SM103 形态）以 MOE_STRATEGY=auto（默认）在单进程内同时执行 prefill 与 decode（PDFUSION，或 prefill 角色进程开启 ENABLE_CUDA_GRAPH），并设 ENABLE_CUDA_GRAPH=1；随后处理长 prompt：默认分块每块 ≤min(moe_prefill_max_tokens_per_rank 或 max_seq_len, 16384) token，DSV4_CHUNK_TOKENS=0 时整段一次进入。
base 中 DSV4 的 MoE 运行时配置 enable_cuda_graph 为类级常量 False（base layer.py:29，构造器无该参数，base deepseek_v4_model.py/transformer.py/block.py 均无引用），graph 开关不进入策略选择，auto 恒选 no_dp hybrid，prefill 大批走 contiguous；no_dp_masked 仅在显式 pin 时可选（base fp8_per_block.py:66）。head 由 379 行注入真实开关并新增 auto_graph_fallback，使 MaskedV2 在 auto 模式自动生效且作用于该层全部前向。MaskedV2 与 execute_masked 的分配代码两版一致，变化仅在选中路径。
head 将引擎 graph 开关贯通到 DSV4 MoE 层（deepseek_v4_model.py:379 → transformer.py:191 → block.py:179 → chunked_layer.py:226 → layer.py:136），并在 fp8_per_block.py:68-78 新增 auto_graph_fallback：auto+graph+ep==1+tp==1+9<=sm<12 时，no_dp 的 hybrid 执行器被 check_conditions(204) 拒绝，DeepEP 两路由要求 ep>1（deepep_normal_router.py:48、deepep_low_latency_router.py:54），pure_dp/pure_cp 仅显式名可选，唯一候选为 no_dp_masked（test_cuda_strategies.py:384-393 断言 SM9/SM10 选中 masked）。DeepGemmMaskedExecutorV2.execute 无条件调用 execute_masked（60 行），后者以 alignment=align(token_num,128) 分配 [E,align,K]、[E,align,N]、[E,align,N//2]、[E,align,K] 等缓冲（493/502/548/565/621 行），全链路不引用 masked_max_token_num。策略在层构造时一次确定（layer.py:267），prefill（block.py:394/578/743）与 decode（479）共用同一 self.ffn；hybrid 执行器 114 行注释明确 enable_cuda_graph 对普通 eager prefill 前向同样为真。分块仅把每次调用 token 数收敛到 min(prefill 预算,16384)（chunked_layer.py:87/278），DSV4_CHUNK_TOKENS=0 时整段不分块，因此数千至 16384 token 的 prefill 同样进入 E×T 级 masked 布局。
ep==1 时 num_experts_per_partition 为全部路由专家数（V4Args 默认 256）。按默认维度（dim=4096、moe_inter_dim=2048 → K=4096、N=4096）估算，execute_masked 并发峰值 T=8192 约 26GB、T=16384 约 52GB 临时缓冲，随 T 线性放大；文档记录的 1×SM103 单卡部署在 KV 分配后已几无空闲显存（曾差 256MiB），该组合一旦出现，首个长 prefill 在 torch.empty 处 CUDA OOM，前向失败、服务中断（引擎对前向异常的终止语义未在快照内定位）；base 同配置 prefill 走 contiguous 不会发生。若空闲显存充足，则退化为工作区数十倍膨胀与显著变慢，不必然崩溃。
rtp_llm/models_py/model_desc/deepseek_v4_model.py:379：args.enable_cuda_graph = bool(
rtp_llm/models_py/modules/dsv4/block.py:179：enable_cuda_graph=enable_cuda_graph,
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:136：self.enable_cuda_graph = bool(enable_cuda_graph)
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:29：enable_cuda_graph = False
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_per_block.py:70：and config.enable_cuda_graph
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_per_block.py:73：and 9 <= get_sm()[0] < 12
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_per_block.py:66：checker.check(config.moe_strategy == "fp8_per_block_no_dp_masked")
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_hybrid_executor.py:204：checker.check(not config.enable_cuda_graph or get_sm()[0] == 12)
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/deepep_normal_router.py:48：checker.check(resolver.is_ep_enabled(config))
rtp_llm/models_py/modules/factory/fused_moe/tests/test_cuda_strategies.py:385：(9, CudaFp8PerBlockNoDPMaskedStrategy),
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_masked_executor_v2.py:60：return self.execute_masked(
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_hybrid_executor.py:493：alignment = align(token_num, self.EXPERT_ALIGNMENT)
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_hybrid_executor.py:502：(self.num_experts_per_partition, alignment, self.K),
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_hybrid_executor.py:548：(self.num_experts_per_partition, alignment, self.N),
rtp_llm/server/server_args/moe_group_args.py:180：help="非deepep low latency场景下使用deepgemm masked的最大token数目, 默认为256",
rtp_llm/models_py/modules/dsv4/block.py:578：ffn_out = self.ffn(x_pre, input_ids)
rtp_llm/models_py/modules/dsv4/block.py:479：ffn_out = self.ffn(x_pre, input_ids, is_decode_forward=True)
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:267：self.fused_moe = FusedMoeFactory().create_fused_moe(
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:87：return max(min(budget, chunk_tokens) if chunking_enabled else budget, 1)
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:278：return self.chunking_enabled and tokens > max_tokens
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_hybrid_executor.py:114：enable_cuda_graph is an engine-level capability and is also true for
rtp_llm/cpp/models/PyWrappedModel.h:346：init_resources.is_decode_role = (params.parallelism_config.role_type == RoleType::DECODE);
rtp_llm/model_loader/per_block_fp8_quant_weight.py:1285：# OOM at KV-cache alloc (256 MiB short on a 267 GiB device with
rtp_llm/models_py/modules/dsv4/transformer.py:80：n_routed_experts: int = 256
rtp_llm/server/server_args/hw_kernel_group_args.py:125：default=False,
待确认：是否存在 ep==1/tp==1 的 FP8-per-block DSV4 部署在开启 ENABLE_CUDA_GRAPH 的同一进程内执行长 prompt prefill（PDFUSION 混合角色或 prefill 角色带 graph 开关）；快照内开 graph 的 DSV4 用例均为 PD 分离、prefill 侧 graph=0 且 pin mega_moe_se，无法证实或证伪外部部署形态; 该部署实际检查点的 E/K/N 维度与运行期空闲显存（决定 E×T 临时缓冲峰值是否必然超出）
建议：把 is_decode_role 传入 Fp8Fp4MoeRuntimeConfig（该标志已在 block.py:181→chunked_layer.py:186 链路可得）并将 auto_graph_fallback 限定于 decode 角色实例；或在 DeepGemmMaskedExecutorV2.execute 中对 token_num 超过 masked_max_token_num 且非 warmup/捕获态的前向回落 execute_contiguous 或显式报错，禁止 prefill 大批按 E×T 无上限分配。
评审版本：f8c79ba72870
| ) | ||
| routed_dtype = layer_weights[W.moe_w1].dtype | ||
| if routed_dtype in {torch.float8_e4m3fn, torch.float8_e4m3fnuz}: | ||
| moe_quant_method = "FP8_PER_BLOCK" |
There was a problem hiding this comment.
[P2] DSV4 FP8 路由权重 + EP>1 + enable_cuda_graph（sm≠12）建层期无可选 MoE 策略，诊断不指向图开关
raw HF 叠放 BF16 DSV4 检查点 + --quantization FP8_PER_BLOCK（LoadQuantPerBlockFp8Weight 在线量化 moe_w1/moe_w2 为 float8_e4m3fn）或预量化 V4 FP8 检查点，多卡 ep_size>1，PDFUSION 或 PD prefill 角色（自动配置 use_deepep_low_latency=False），--enable_cuda_graph 1，sm90/sm100；deepseek_v4_model→V4Transformer→Block→ChunkedFp8Fp4MoeLayer→Fp8Fp4MoeLayer→create_fused_moe 建层即触发。
base 的 Fp8Fp4MoeRuntimeConfig 恒为类属性 moe_quant_method=FP8_FP4、enable_cuda_graph=False，DSV4 描述符只有 int8 打包 FP4 路由专家（mega 族无图检查），Block 不透传图开关，该组合不可达；base 虽已有 hybrid 执行器与 pure_dp 路由器的图守卫，但对 DSV4 是死代码。head 的 dtype 推导、四级透传与 raw BF16 在线量化路径首次使 FP8+EP>1+graph 组合到达这些守卫，而图兼容 EP 后端仅覆盖 sm120 与 decode 低时延。
head 按路由权重 dtype 推导 moe_quant_method=FP8_PER_BLOCK 并构造 Fp8BlockWiseQuantConfig 后，fp8_fp4/mega 族被注册表 supported_moe_quant_method 过滤与策略量化条件双重排除，per_tensor/fp4/no_quant/w4a8/batched_triton 族被各自执行器或策略的量化检查拒绝，仅 fp8_per_block 六策略可候选；can_handle 串联策略/路由器/执行器三级检查：graph 开启且 ep>1、sm≠12、moe_strategy=auto、use_deepep_low_latency=False 时，no_dp 与 ep_normal 被 DeepGemmHybridExecutor 的 not enable_cuda_graph or sm==12 拒绝，no_dp_masked 的 auto 回退要求 ep_size==1 且 tp_size==1，pure_dp/pure_cp 需显式策略名（pure_dp 路由器另拒图），ep_low_latency 路由器要求 use_low_latency——候选为空，注册表抛 ValueError，文案与 error 日志仅列 quant_config/ep_size/world_size/tp_size/use_deepep_low_latency，不含 enable_cuda_graph/SM 维度。
模型初始化建层期抛 ValueError，fail-fast 响亮失败、无静默数值错误；但报错文案与日志引导用户排查 quant_config/ep_size/parallelism 而非图开关/SM 版本，排障成本高；FP8 路由权重 + EP>1 部署在 sm90/100 上被迫关闭 CUDA graph 才能启动。属受支持非默认路径的启动可用性缺口与诊断缺陷，非数据错误。
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:232：moe_quant_method = "FP8_PER_BLOCK"
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:238：f"got {routed_dtype}. Raw BF16 checkpoints require "
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_fp4.py:30：checker.check(config.moe_quant_method == "FP8_FP4")
rtp_llm/models_py/modules/factory/fused_moe/strategy_registry.py:174："No suitable MOE strategy found for configuration. "
rtp_llm/models_py/modules/factory/fused_moe/strategy_registry.py:175："Please check quant_config, ep_size, and parallelism settings. "
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_hybrid_executor.py:204：checker.check(not config.enable_cuda_graph or get_sm()[0] == 12)
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_per_block.py:71：and config.ep_size == 1
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_per_block.py:118：checker.check(config.moe_strategy == "fp8_per_block_pure_dp")
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_per_block.py:161：checker.check(config.moe_strategy == "fp8_per_block_pure_cp")
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_dp_router.py:85：checker.check(not config.enable_cuda_graph)
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/deepep_low_latency_router.py:55：checker.check(resolver.use_low_latency(config))
rtp_llm/models_py/model_desc/deepseek_v4_model.py:380：getattr(py_hw_kernel_config, "enable_cuda_graph", False)
rtp_llm/config/server_config_setup.py:206：use_deepep_low_latency = False
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:27：moe_quant_method = "FP8_FP4"
rtp_llm/models/deepseek_v4.py:427：data_type=torch.int8,
建议：在注册表失败文案与 error 日志中加入 enable_cuda_graph/SM 维度；DSV4 建层前对 FP8_PER_BLOCK+ep>1+graph+sm≠12 组合给出明确启动期指引（关闭图、走 sm120 路径或 decode 角色启用 DeepEP 低时延）；或在验证后为 sm9x/10x 提供图兼容的 EP 执行器。
评审版本：f8c79ba72870
| (0, v4_args.dim), dtype=v4.embed.weight.dtype, device=param_dev | ||
| ) | ||
| empty_ids = torch.empty((0,), dtype=torch.long, device=param_dev) | ||
| with synchronized_moe_chunk_plan(v4.layers, 0, param_dev): |
There was a problem hiding this comment.
[P2] pure_dp 全 rank 同时空批的解码步以 grid=(0,) 启动 recompute_topk_ids_sum_expert_count（零输入无早退）
显式配置 moe_strategy=fp8_per_block_pure_dp（tp_size==1、dp_size>1、ep_size==dp_size、eager 非图路径），且引擎在某步以空模型输入向全部 DP rank 发起解码 forward。快照内唯一可见的空输入前向引擎来源是 ffn-disaggregate 对齐（NormalExecutor.cc:268 skip_run 为假；MtpExecutor alignment_only → runDecodeAlignmentOnly 以空输入调用 forwardModel）；正常引擎 dp>1 时空闲 rank 被 1-token fake stream 填充（NormalEngine.cc:985），全空闲步以 bs≥1 前向，不进入该空分支。
base 在 meta is None 处直接短路返回零行输出、不调用任何 ffn（base decode/forward.py:469-472），且 base 同步计划仅覆盖 mega_moe/mega_moe_se（base chunked_layer.py:108）；head 将 pure_dp/ep_normal 纳入集体策略集合并新增空批 0 行参与路径，使 base 已存在但空解码路径不可达的零网格启动（ep_kernels.py、pure_dp_router.py 两文件本身未变更）成为新引入的可达路径。
head 空批参与路径在 meta is None 时对 strategy_name ∈ _COLLECTIVE_CHUNK_STRATEGIES（含 fp8_per_block_pure_dp）的层以 0 行调用 ffn；ChunkedFp8Fp4MoeLayer 对 0 行无早退（全空时同步计划 local_chunk_count=max(...,1)=1、planned_chunks=1，单次 call_moe），FusedMoeDataRouter.forward 无条件调用 router.prepare；PureDpRouterBase.prepare 中 pad_to_max 在全 rank 空批时 max_n=0、不填充，all_gather 后 topk_ids_full 为 [0,k]，recompute_topk_ids_sum_expert_count 得 num_total=0、grid=(cdiv(0,BLOCK),)=(0,) 并直接启动 triton 内核。0 行路径上 gate（0 行走 eager，fused gate N==0 早退）、norm_top_p（纯 torch）、shared expert（T==0 早退）、量化（x.shape[0]>0 守护）均有守护，仓内代码未守护点仅剩 0 行 all_gather_into_tensor 与该零网格启动。
若 0 行 all_gather 可通过且所用 triton 对 grid=(0,) 启动报错，则每个全空解码步在每个 pure_dp MoE 层抛错、整步失败，并可能在空闲/对齐期反复发生；若 0 行 all_gather 先报错，则同一链路更早位置发生同性质的步失败；若外部零规模行为均放行则无影响。定级按未知前提成立时的后果（受支持非默认路径运行时崩溃）计为 P1，非置信度。
rtp_llm/models_py/modules/dsv4/decode/forward.py:484： with synchronized_moe_chunk_plan(v4.layers, 0, param_dev):
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:33： "fp8_per_block_pure_dp",
rtp_llm/models_py/modules/factory/fused_moe/defs/fused_moe.py:341： expert_payload = self.router.prepare(
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_dp_router.py:123： max_n = int(all_n.max().item())
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_dp_router.py:152： recompute_topk_ids_sum_expert_count(
rtp_llm/models_py/triton_kernels/moe/ep_kernels.py:646： num_total = num_tokens * topk
rtp_llm/models_py/triton_kernels/moe/ep_kernels.py:659： grid_recompute = (triton.cdiv(num_total, BLOCK_SIZE),)
rtp_llm/models_py/triton_kernels/moe/gate_fused.py:186： if N == 0:
rtp_llm/models_py/modules/dsv4/decode/forward.py:469： # Empty batch (B == 0) — short-circuit with zero-row hidden.
rtp_llm/cpp/normal_engine/NormalEngine.cc:985： streams.emplace_back(createMinFakeStream(1));
rtp_llm/cpp/normal_engine/NormalExecutor.cc:268： model_input.skip_run = streams.empty() && !enable_ffn_disaggregate;
rtp_llm/models_py/modules/dsv4/fp8/test/test_topk_v3.py:779： """A zero-row tensor must return without launching a zero-sized grid."""
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_per_block.py:118： checker.check(config.moe_strategy == "fp8_per_block_pure_dp")
rtp_llm/models_py/distributed/collective_torch.py:847： torch.distributed.all_gather_into_tensor(tensor_list, tensor, group=process_group)
rtp_llm/cpp/normal_engine/speculative/MtpExecutor.cc:2076： forwardModel(model.get(), model_input, ModelInputsModelRole::TARGET);
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/shared_expert.py:407： if T == 0:
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/gate.py:61： if x_size_0 == 0 or not 1 <= topk <= 32:
deps/requirements_torch_gpu_cuda13.txt:17：triton==3.6.0
rtp_llm/models_py/kernels/cuda/fp8_kernel/fp8_kernel.py:135： if x.shape[0] > 0:
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:108： in {"mega_moe", "mega_moe_se"}
rtp_llm/models_py/triton_kernels/moe/ep_kernels.py:659： grid_recompute = (triton.cdiv(num_total, BLOCK_SIZE),)
待确认：pure_dp（DSV4）部署中是否存在全部 DP rank 同为 0 token 的解码步被发起，含 ffn-disaggregate 与 DSV4/pure_dp 组合是否受支持（快照内唯一空输入前向来源即该对齐路径）; torch/NCCL 的 all_gather_into_tensor 对 0 行张量的行为（放行还是抛错）; 所锁定 triton 版本（3.2.0–3.8.0）对 grid=(0,) 内核启动的行为（报错还是跳过）
建议：在 recompute_topk_ids_sum_expert_count 开头对 num_tokens==0（num_total==0）直接返回（同形空 adjusted_topk_ids、全零 expert_count），或在 PureDpRouterBase.prepare 中对 max_n==0 短路返回空 payload，与 gate_fused/shared_expert/mega_moe_input_pack 的既有零规模早退惯例一致；同时对空批参与路径上的 0 行集合通信补充守护或验证。
评审版本：f8c79ba72870
| # as the invalid-slot sentinel. The unquantized fused-MoE executor | ||
| # indexes the fixed-width IDs directly, so it still needs expert zero | ||
| # as a placeholder for deleted slots. | ||
| valid_ids = (topk_ids >= 0) & (topk_ids < self.expert_num) |
There was a problem hiding this comment.
[P2] PureTpRouterBase.prepare 在 do_recompute_topk 分支前无条件构造三份净化张量：每次 CUDA 纯 TP MoE 前向新增约 9 次 eager 内核与 7 个新张量，且每个分支各含约 2 次死操作
任意 CUDA 纯 TP MoE 前向都会执行：bf16 默认部署（moe_strategy=no_quant_cpp/auto → PureTpRouterNoQuant，do_recompute_topk=False，no_quant.py:77、pure_tp_router.py:166）、DSV4 FP8 部署（fp8_per_block_no_dp(_masked) → PureTpRouterFp8PerBlock，True，fp8_per_block.py:50、pure_tp_router.py:218）、以及 FP8 per-tensor/W4A8/FP4 等继承同一 prepare 的量化路由；包括完全未启用 top-p 的部署。
base 直接把原始 topk_ids 交给重算内核（110/116 行）并把调用方 topk_weights 原样放进 payload（125 行）；head 每次前向额外执行上述约 9 次内核与 7 个张量分配，payload 权重变为新分配的 safe_topk_weights，与 top-p 是否启用、与 do_recompute_topk 取值均无关。
head 在 if self.do_recompute_topk（128 行）之前无条件计算 valid_ids（2 次比较 + 1 次逻辑与）、safe_topk_ids（zeros_like + where，钳 0）、recompute_topk_ids（full_like(-1) + where）、safe_topk_weights（zeros_like + where，置 0），eager 下共约 9 次设备内核启动，新增 3 个输出张量与 4 个掩码/填充临时张量，无任何 top-p 或配置门控；生产代码 torch.compile 仅包住 moe_align_block_size_torch（fused_moe_kernel.py:372，全仓检索无其他生产调用），不覆盖路由器 prepare。do_recompute_topk=True 时 adjusted_topk_ids 先被赋 safe_topk_ids（125 行）随后被内核返回值整体覆盖（129-135 行），safe_topk_ids 的 2 次操作是死代码；=False 时 recompute_topk_ids 的唯一引用在 if 分支内（131 行），其 2 次操作同样是死代码。
每 MoE 层每次前向固定新增约 9 次内核启动与 7 个张量分配，其中约 2 次操作在当前分支下必为死代码：eager 解码下随层数线性放大为额外启动时延，CUDA graph 下为图内额外小内核执行与内存池占用；payload 权重从复用调用方张量变为每步新分配。所有可达支持配置下数值与 base 等价，无正确性影响；量级未实测，属默认热路径效率缺陷。
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_tp_router.py:115： valid_ids = (topk_ids >= 0) & (topk_ids < self.expert_num)
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_tp_router.py:117： valid_ids, topk_ids, torch.zeros_like(topk_ids)
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_tp_router.py:119： recompute_topk_ids = torch.where(
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_tp_router.py:120： valid_ids, topk_ids, torch.full_like(topk_ids, -1)
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_tp_router.py:122： safe_topk_weights = torch.where(
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_tp_router.py:125： adjusted_topk_ids = safe_topk_ids
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_tp_router.py:128： if self.do_recompute_topk:
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_tp_router.py:131： recompute_topk_ids,
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_tp_router.py:142： safe_topk_weights,
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_tp_router.py:110： adjusted_topk_ids = topk_ids
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_tp_router.py:116： topk_ids, self.expert_start_id, self.expert_num_per_rank
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_tp_router.py:125： topk_weights,
rtp_llm/models_py/triton_kernels/moe/ep_kernels.py:614： valid = mask & (adjusted >= 0) & (adjusted < num_local_experts)
rtp_llm/models_py/triton_kernels/moe/ep_kernels.py:617： out = tl.where(valid, adjusted, -1)
rtp_llm/models_py/triton_kernels/moe/ep_kernels.py:621： tl.atomic_add(expert_count_ptr + adjusted, 1, mask=valid)
rtp_llm/models_py/triton_kernels/moe/ep_kernels.py:617： out = tl.where(valid, adjusted, -1)
rtp_llm/models_py/modules/dsv4/norm_predictor.py:188： kept_weights = topk_weights.masked_fill(~keep, 0)
rtp_llm/models_py/modules/dsv4/norm_predictor.py:260： invalid_expert_id=self.n_routed_experts,
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:271： if self.norm_top_p.enabled and self.strategy_name not in {
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:237： "DeepSeek-V4 routed experts must be FP8-per-block or packed FP4; "
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:266： runtime_config.norm_top_p_enabled = self.norm_top_p.enabled
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_per_block.py:50： router_class=PureTpRouterFp8PerBlock,
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_tp_router.py:218： super().init(config, quant_config, do_recompute_topk=True)
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/no_quant.py:77： router_class=PureTpRouterNoQuant,
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_tp_router.py:166： super().init(config, quant_config, do_recompute_topk=False)
rtp_llm/models_py/modules/factory/fused_moe/defs/fused_moe.py:341： expert_payload = self.router.prepare(
rtp_llm/models_py/triton_kernels/moe/fused_moe_kernel.py:294： Requires every route id to be in [0, num_experts). EP padding
rtp_llm/models_py/triton_kernels/moe/fused_moe_kernel.py:372：moe_align_block_size_compiled = torch.compile(
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/triton_fused_executor.py:88： topk_ids = local_ids.clamp(min=0, max=self.num_experts_per_partition - 1)
rtp_llm/models_py/modules/factory/fused_moe/impl/rocm/routers/pure_tp_router.py:117： topk_ids, self.expert_start_id, self.expert_num_per_rank
rtp_llm/models_py/model_desc/generic_moe.py:211： self.select_topk(router_logits, topk_ids, topk_weights)
rtp_llm/models_py/bindings/cuda/SelectTopkOp.cc:36： num_expert,
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/deepep_normal_router.py:115： valid_ids = (topk_ids >= 0) & (topk_ids < self.expert_num)
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:223： self.norm_top_p = Dsv4NormTopP(
建议：把净化下放到分支并按 top-p 门控：do_recompute_topk=True 时直接传原始 topk_ids（内核已折叠越界为 -1），仅在需要时构造 safe_topk_weights；do_recompute_topk=False 时仅对确可收到越界 id 的部署构造 expert-0 占位。可复用 layer.py:266 已写入的 runtime_config.norm_top_p_enabled（deepgemm_hybrid_executor.py:232 已用同名属性做门控）使未启用 top-p 的默认部署整体跳过；若确需统一防御，用单个融合算子一次产出两份结果，消除约 9 次启动、7 个张量与各分支约 2 次死操作。deepep 路由器中净化被实际消费（deepep_normal_router.py:124），应保留、按路由器区分修复。
评审版本：f8c79ba72870
| ) | ||
| self.strategy_name = self.fused_moe.strategy_name | ||
| if self.norm_top_p.enabled and self.strategy_name not in { | ||
| "fp8_per_block_no_dp", |
There was a problem hiding this comment.
[P2] Top-p 白名单仅按策略名放行 fp8_per_block_no_dp，head 新增的 fnuz/fp8→FP8_PER_BLOCK 解析使 ROCm 同名策略首次可从该层选中，其透传路由把哨兵专家 id(=n_routed_experts) 未消毒直送 aiter fused_moe
在可构建 DSV4 models_py 栈的 ROCm 部署上以 FP8-per-block 权重（加载期转 fnuz）服务 deepseek_v4，纯 TP（ep_size=1 且 dp_size=1、use_all_gather），设置 RTP_DSV4_NORM_PREDICTOR_DIR 与 RTP_DSV4_NORM_TOP_P<1.0，策略 auto 或显式解析为 fp8_per_block_no_dp；layer_id>=n_hash_layers 的任一 token 出现被剪枝槽即触发。
base 无 norm_predictor 与白名单，层恒置 moe_quant_method=FP8_FP4 且 quant_config=None，ROCm FP8-per-block 策略的 check_conditions 不可能通过，任何路由也收不到越界 id；head 新增哨兵生产者、按名白名单与 fnuz/fp8→FP8_PER_BLOCK 解析，使 ROCm 同名策略 newly 可达白名单，但仅为 CUDA 路由/执行器加了消毒，ROCm 侧未同步。
Dsv4NormTopP 把被剪枝槽写为 id=n_routed_experts（DSV4 为 256）且权重 0；Fp8Fp4MoeLayer 的 Top-p 门禁只比较 strategy_name 字符串并放行 "fp8_per_block_no_dp"，且 head 新增 dtype 分支把 e4m3fn/fnuz 权重解析为 FP8_PER_BLOCK 并接入 Fp8BlockWiseQuantConfig；ROCm 分支注册表将同名策略映射到 RocmFp8PerBlockPureTPStrategy（策略级 check_conditions 为基类 no-op），其 PureTpRouterFp8PerBlockPassthrough 为 do_recompute_topk=False 且 prepare 原样透传 topk_ids/topk_weights，RocmExpertsFp8PerBlock 在 pure TP 下 expert_mask=None，把含 id=256 的 topk_ids 直接传入外部 aiter fused_moe；该执行器也不读取 head 新设的 norm_top_p_enabled。
若两个未知前提成立：每个受影响 MoE 层的被删槽以 id=expert_num 越出 [0, expert_num) 进入外部 aiter 内核，可能设备端越界索引崩溃，或经越界读产生 NaN/Inf 污染输出（槽权重为 0 不能一般化消除 NaN 传播）；若 aiter 钳制或安全跳过越界 id，则仅剩契约违规无数值后果。属 env 门控+ROCm 的受支持非默认路径上的崩溃/数据损坏类失败。
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:272："fp8_per_block_no_dp",
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:281："that skips dropped route slots; got "
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:231：if routed_dtype in {torch.float8_e4m3fn, torch.float8_e4m3fnuz}:
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:232：moe_quant_method = "FP8_PER_BLOCK"
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:141：Fp8BlockWiseQuantConfig(is_quanted=True)
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:373：topk_weights, topk_ids = self.norm_top_p(
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:389：return self.fused_moe(
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:27：moe_quant_method = "FP8_FP4"
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:113：self.quant_config = None
rtp_llm/models_py/modules/dsv4/norm_predictor.py:190：kept_ids = torch.where(keep, topk_ids, invalid)
rtp_llm/models_py/modules/dsv4/norm_predictor.py:260：invalid_expert_id=self.n_routed_experts,
rtp_llm/models_py/modules/dsv4/norm_predictor.py:188：kept_weights = topk_weights.masked_fill(~keep, 0)
rtp_llm/models_py/modules/dsv4/norm_predictor.py:24：INVALID_EXPERT_ID = 256
rtp_llm/models_py/modules/dsv4/norm_predictor.py:226：and layer_id >= n_hash_layers
rtp_llm/models_py/modules/factory/fused_moe/__init__.py:66：registry.register(RocmFp8PerBlockPureTPStrategy())
rtp_llm/models_py/modules/factory/fused_moe/impl/rocm/strategy/fp8_per_block.py:18：strategy_name = "fp8_per_block_no_dp"
rtp_llm/models_py/modules/factory/fused_moe/impl/rocm/routers/pure_tp_router.py:111：adjusted_topk_ids = topk_ids
rtp_llm/models_py/modules/factory/fused_moe/impl/rocm/routers/pure_tp_router.py:202：super().init(config, quant_config, do_recompute_topk=False)
rtp_llm/models_py/modules/factory/fused_moe/impl/rocm/routers/pure_tp_router.py:209：checker.check(quant_method in ("FP8_PER_BLOCK", "FP8_PER_BLOCK_QUARK"))
rtp_llm/models_py/modules/factory/fused_moe/impl/rocm/executors/rocm_moe.py:6：from aiter.fused_moe import fused_moe
rtp_llm/models_py/modules/factory/fused_moe/impl/rocm/executors/rocm_moe.py:371：output = fused_moe(
rtp_llm/models_py/modules/factory/fused_moe/impl/rocm/executors/rocm_moe.py:45：if ep_size <= 1:
rtp_llm/models_py/modules/factory/fused_moe/defs/strategy_base.py:101：pass
rtp_llm/config/quant_config.py:550：return "FP8_PER_BLOCK"
rtp_llm/models_py/modules/factory/fused_moe/strategy_registry.py:28：return quant_method != "FP8_FP4" or not isinstance(strategy, MoeStrategy)
rtp_llm/models_py/modules/factory/fused_moe/utils/config_resolver.py:139：return config.ep_size == 1 and config.dp_size == 1 and config.tp_size >= 1
rtp_llm/models_py/modules/factory/fused_moe/defs/fused_moe.py:341：expert_payload = self.router.prepare(
rtp_llm/models_py/modules/factory/fused_moe/impl/common/router/batched_data_router.py:50：checker.check(not resolver.has_quantization(config))
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_tp_router.py:115：valid_ids = (topk_ids >= 0) & (topk_ids < self.expert_num)
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_tp_router.py:110：# Keep the fixed-width Top-p contract safe for both router families.
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_tp_router.py:110：adjusted_topk_ids = topk_ids
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/test/pure_tp_router_skip_allreduce_test.py:67：topk_ids = torch.tensor([[3, 256], [1, -1]], dtype=torch.int32)
rtp_llm/models_py/modules/factory/fused_moe/utils/test/fp8_fp4_layer_test.py:187："fp8_per_block_no_dp",
rtp_llm/models_py/modules/dsv4/fp8/attention.py:28：import deep_gemm # noqa: E402
rtp_llm/models_py/modules/dsv4/fp8/attention.py:1942：# absent (e.g. ppu/rocm), so importing it must not require flashinfer.
rtp_llm/models_py/modules/dsv4/fp8/attention.py:28：import deep_gemm # noqa: E402
deps/requirements_rocm.txt:8：https://sinian-metrics-platform.oss-cn-hangzhou.aliyuncs.com/kis/AMD/aiter/amd_aiter-0.1.23.dev134%2Bg8449f416b-cp310-cp310-linux_x86_64.whl
3rdparty/deep_gemm/BUILD:54："-DWITH_CUDA",
rtp_llm/device/device_impl.py:963：weight = weight_as_int8.view(torch.float8_e4m3fnuz)
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:266：runtime_config.norm_top_p_enabled = self.norm_top_p.enabled
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_hybrid_executor.py:232：self.norm_top_p_enabled = bool(getattr(config, "norm_top_p_enabled", False))
待确认：aiter fused_moe 对 topk_id==expert_num（越界）的实际处理（崩溃、静默误路由或安全跳过）——外部 wheel 源码不在快照内，决定失败是否存在; 目标 ROCm 部署是否提供可导入的 deep_gemm（或等效物）使 DSV4 model_desc 构建成功——决定场景可达性；仓库 ROCm requirements/lock 均无 deep_gemm，而 attention.py:1941-1942 注释与 head 新增 fnuz 支持指向相反方向
建议：白名单改按路由/执行器实际消毒能力或平台判定（放行前校验该策略的 router 会跳过无效槽）；或在 impl/rocm/routers/pure_tp_router.py 的 prepare 加入与 CUDA 一致的 valid_ids 消毒（0 占位+零权重或 -1 哨兵）；或在非 CUDA 设备上禁止 Dsv4NormTopP 启用并在 layer 守卫显式报错。
评审版本：f8c79ba72870
| and config.enable_cuda_graph | ||
| and config.ep_size == 1 | ||
| and config.tp_size == 1 | ||
| and 9 <= get_sm()[0] < 12 |
There was a problem hiding this comment.
[P2] auto+CUDA Graph 兜底使单卡 FP8-per-block MoE 的 eager 预填充进入无 token 上界的 E×align(T,128) 掩码布局，与默认约 5% 运行期显存余量冲突，服务期可致 CUDA OOM
树内已存在的单卡拓扑（H20 smoke：qwen_3_moe + FP8_PER_BLOCK + tp=1/dp=1/ep=1 + enable_cuda_graph 1）保持 moe_strategy 默认 auto、reserver_runtime_mem_mb 默认 1024、warm_up 默认 True；启动 warmup 以 max_seq_len−1（qwen_3_moe 默认 max_seq_len=8192）在 KV 分配前通过，随后 KV 吞掉除约 5% 地板外的显存；服务期一条接近 max_seq_len 的整长预填充进入 execute_masked，峰值约 5GiB。
base 中 masked 策略仅接受显式策略名（base fp8_per_block.py:66），auto+enable_cuda_graph+SM9-11+tp=ep=1 无任何候选（hybrid 的 not enable_cuda_graph or sm==12 检查 base/head 相同，deepep 两路由均要求 ep>1），注册表抛 ValueError 启动即失败；head 新增 auto_graph_fallback（fp8_per_block.py:68-78）后同一配置可运行且唯一选中 V2，作者测试即断言 SM9/SM10 的 auto+图选中 masked 策略，但新开通路径未附带任何 token/显存预算约束。
head 的 auto_graph_fallback 只读引擎级静态开关 enable_cuda_graph（generic_moe.py:88 在建层时固化，对普通 eager 预填充同样为真），把默认 auto 的单卡 SM9–11 FP8_PER_BLOCK 配置唯一路由到 DeepGemmMaskedExecutorV2；V2.execute 无条件调用继承的 execute_masked，按 num_experts_per_partition×align(token_num,128) 分配 input(fp8,K)+scale+upgate(bf16,N)+down_input(fp8,N/2)+down_output(bf16,K)，峰值约 0.63MiB/token（E=128、K=2048、N=1536），完全不检查 masked_max_token_num（该阈值仅存在于 hybrid 的 token_num<=masked_max_token_num 且非图准入）。GenericMoeLayer.forward 整批直通不分块，调度侧 max_batch_tokens_without_cache 默认 0（不限），max_context_batch_size 默认 1 但长度可达 max_seq_len。auto 路径不满足生成期预填充图的显式策略名条件（仅记 warning 后禁用），故预填充以整批 token 走掩码布局。引擎侧 KV 按 available − runtime_required 预分配，runtime_required = max(reserve_runtime_mem_mb 默认 1024, max(2GiB, 总显存 5%), warmup 峰值)，而快照内唯一 getGpuExecStatus 只填 free/used/available，max_consumed_bytes 恒为初值 0，故 warmup（长度 max_seq_len−1，在 KV 分配前、显存充裕时执行）测得的掩码峰值不会扩大预留，服务期可用余量停留在约 5% 地板。
预填充瞬态超出运行期余量（总显存低于约 101GiB 时 5% 地板即小于该 5GiB 峰值）→ torch CUDA OOM；引擎 step() 仅在转储诊断后重抛，loop() 无捕获，异常逃逸引擎循环线程，长上下文请求持续失败并可致服务不可用；base 同配置为启动期干净 ValueError，head 把失败点推迟到服务期。未越阈时退化为相对 contiguous（约 top_k×T 行）约 E/top_k≈16 倍的行数/显存放大并挤占 KV 容量。
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_per_block.py:73：and 9 <= get_sm()[0] < 12
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_per_block.py:77：or auto_graph_fallback
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_per_block.py:66：checker.check(config.moe_strategy == "fp8_per_block_no_dp_masked")
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_hybrid_executor.py:204：checker.check(not config.enable_cuda_graph or get_sm()[0] == 12)
rtp_llm/models_py/modules/factory/fused_moe/strategy_registry.py:174："No suitable MOE strategy found for configuration. "
rtp_llm/models_py/modules/factory/fused_moe/tests/test_cuda_strategies.py:385：(9, CudaFp8PerBlockNoDPMaskedStrategy),
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/deepep_low_latency_router.py:54：checker.check(resolver.is_ep_enabled(config))
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_tp_router.py:66：checker.check(resolver.use_all_gather(config))
rtp_llm/server/server_args/moe_group_args.py:198：default="auto",
rtp_llm/server/server_args/moe_group_args.py:187：default=True,
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_masked_executor_v2.py:60：return self.execute_masked(
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_hybrid_executor.py:326：if token_num <= self.masked_max_token_num and not self.enable_cuda_graph:
rtp_llm/server/server_args/moe_group_args.py:180：help="非deepep low latency场景下使用deepgemm masked的最大token数目, 默认为256",
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_hybrid_executor.py:493：alignment = align(token_num, self.EXPERT_ALIGNMENT)
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_hybrid_executor.py:502：(self.num_experts_per_partition, alignment, self.K),
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_hybrid_executor.py:548：(self.num_experts_per_partition, alignment, self.N),
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_hybrid_executor.py:620：down_output = torch.empty(
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_hybrid_executor.py:114：enable_cuda_graph is an engine-level capability and is also true for
rtp_llm/models_py/model_desc/generic_moe.py:88：enable_cuda_graph=enable_cuda_graph,
rtp_llm/models_py/model_desc/generic_moe.py:221：experts_output = self.fused_moe(
rtp_llm/server/server_args/hw_kernel_group_args.py:126：help="系统是否允许使用Cuda Graph",
rtp_llm/cpp/cache/MemoryEvaluationHelper.cc:59：const auto minimal_runtime_bytes = std::max(2048L * 1024 * 1024, (long)(total_gpu_bytes * 0.05));
rtp_llm/cpp/cache/MemoryEvaluationHelper.cc:110：runtime_required_bytes = std::max(env_runtime_required_bytes, warm_up_result->max_used_memory);
rtp_llm/cpp/cache/MemoryEvaluationHelper.cc:136：const auto kv_cache_mem_size = device_reserved_memory_bytes - runtime_required_bytes;
rtp_llm/models_py/bindings/core/ExecOps.cc:585：mem.available_bytes = mem.free_bytes;
rtp_llm/models_py/bindings/core/DeviceData.h:37：size_t max_consumed_bytes = 0; // only applicable if RTP_LLM_TRACE_MEMORY is enabled.
rtp_llm/cpp/normal_engine/NormalEngine.cc:444：const auto max_consumed = getGpuExecStatus().device_memory_status.max_consumed_bytes;
rtp_llm/cpp/engine_base/EngineInitParams.h:28：return max_seq_len - 1;
rtp_llm/server/server_args/device_resource_group_args.py:66：default=1024,
rtp_llm/cpp/models/GenerationPrefillCudaGraphEligibility.h:112：&& moe_config.moe_strategy == "fp8_per_block_no_dp_masked" && moe_config.use_all_gather
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_hybrid_executor.py:72：(128, 1536, 2048, 8): {
rtp_llm/models/qwen_v2.py:352：config.max_seq_len = 8192
rtp_llm/test/smoke/suites_h20_oss.bzl:85：--moe_strategy fp8_per_block_no_dp_masked --quantization FP8_PER_BLOCK --warm_up 0 --act_type BF16 --tp_size 1 --dp_size 1 --ep_size 1
rtp_llm/cpp/normal_engine/NormalEngine.cc:687：auto status = step();
rtp_llm/server/server_args/fifo_scheduler_group_args.py:31：default=0,
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_masked_executor_v2.py:60：return self.execute_masked(
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:18：DEFAULT_MOE_CHUNK_TOKENS = 16384
rtp_llm/models_py/modules/factory/fused_moe/tests/test_cuda_strategies.py:387：(12, CudaFp8PerBlockNoDPStrategy),
rtp_llm/test/smoke/suites_h20_oss.bzl:103：--quantization W4A8_INT4_PER_CHANNEL --warm_up 0 --act_type BF16 --reserver_runtime_mem_mb 16005
rtp_llm/models_py/modules/dsv4/transformer.py:79：moe_inter_dim: int = 2048
rtp_llm/models_py/modules/dsv4/transformer.py:80：n_routed_experts: int = 256
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:232：moe_quant_method = "FP8_PER_BLOCK"
rtp_llm/cpp/normal_engine/NormalEngine.cc:443：THROW_IF_STATUSOR_ERROR(preRun(fake_input, preRunMode::prefill_warm_up));
待确认：部署卡总显存、实际模型形状（E/K/N）与 max_seq_len 及生产请求长度分布：决定 E×align(T,128) 瞬态是否越过约 max(2GiB,5%) 的运行期余量；越阈即 OOM，未越阈仅为显存放大与 KV 挤占
建议：为 auto 兜底路径补上与显式路径等效的规模约束：在 DeepGemmMaskedExecutorV2.execute 中，非图捕获/warmup（可复用 _is_cuda_graph_warmup_or_capture）且 token_num > masked_max_token_num 时，改走按 ≤masked_max_token_num 的 token 切片循环执行掩码内核，或回落继承的 execute_contiguous（后者在 enable_cuda_graph=True 时走 GPU 计数分支，SM90 上该组合当前未被任何路径使用，需先验证 ep_scatter 的 GPU 计数+align_m=128 行为）；同时收紧兜底条件使其仅在图捕获/桶化场景生效，并在命中兜底时校验或提示 reserver_runtime_mem_mb 与 max_seq_len 的显存需求。
评审版本：f8c79ba72870
There was a problem hiding this comment.
[P1] 非 decode 角色显式开 LL 的 dsv4 FP8-per-block 实例启动期全绿（warmup 前向 kv_cache=None 全零短路绕过守卫），首个 prefill 请求的守卫 RuntimeError 沿 step/loop 无捕获传播，引擎服务线程永久终止，实例整体停服并随重启形成崩溃循环
dsv4 FP8-per-block checkpoint + ep_size>1 + moe_strategy=auto + PDFUSION 或 PREFILL 角色实例上显式设置 USE_DEEPEP_LOW_LATENCY=1（帮助文本无角色限制，server_config_setup 直接复制显式值；auto DeepEP 配置仅为 DECODE 角色启用 LL，故需显式误配）。默认 WARM_UP=True、use_batch_decode_scheduler 默认 False 下：构造、模型侧预热、引擎 prefill 预热全部通过，实例带病上线；上线后第一个 prefill 请求即触发。
base 中 dsv4 仅描述 int8-packed FP4 路由权重（base deepseek_v4.py:427，全文无 stacked/float8_e4m3fn 路由路径），Fp8Fp4MoeRuntimeConfig 硬编码 moe_quant_method='FP8_FP4'（base layer.py:27）且内部 moe_config.use_deepep_low_latency=False（base layer.py:121）、不接收外部 moe_config——LL 对 base 的任何 dsv4 部署不可达。head 新增 raw_hf_stacked/build_routed_experts_stacked（head deepseek_v4.py:595）与 dtype→FP8_PER_BLOCK 推导（head layer.py:230-234）、args.moe_config 透传（deepseek_v4_model.py:378），首次使 LL 对 dsv4 可达，并把不兼容暴露点放在运行期 forward 守卫；对 base 可服务的 int8(FP4) checkpoint，head 仍推导 FP8_FP4（234 行）并经注册表过滤维持 mega 选择、行为不变——失败面为本 PR 新引入，而非 base 回归。
auto+显式 USE_DEEPEP_LOW_LATENCY=1+ep>1 时，dsv4 FP8-per-block 路由权重（float8_e4m3fn dtype 推导 FP8_PER_BLOCK）使 LL 策略成为唯一候选：ep_normal 路由因 not use_low_latency 被排除，pure_tp 族因显式 LL 使 use_all_gather 置 False（server_config_setup.py:107-112）而被 pure_tp_router.py:66 排除，mega 族因 FP8_FP4 量化方法限定（fp8_fp4.py:30）与注册表过滤（strategy_registry.py:28）被排除；LL 策略 check_conditions（fp8_per_block.py:189-196）与 LL 路由检查均不感知部署角色（is_decode_role 只存在于 C++ 注入的 init_resources，不进入 MoE 配置）。唯一的 prefill 拒绝守卫在 ChunkedFp8Fp4MoeLayer.forward:348-354，构造不触发；模型侧 MoE 预跑仅限 mega 策略（deepseek_v4_model.py:931-934）；引擎默认 WARM_UP=True 且 use_batch_decode_scheduler 默认 False 时 PDFUSION/PREFILL 走 prefillWarmUp，虽发出真实 prefill 前向，但 warmup executor 以 cache_manager=nullptr 构造（NormalEngine.cc:442）→kv layout 置 nullopt（NormalExecutor.cc:127-130）→PyWrappedModel.h:341 门控不注入 kv_cache→Python 侧 kv_cache=None→dsv4 forward 短路返回全零（1370-1381）、不进任何层，守卫不被触碰，启动全绿。服务期生产 executor 绑定真实 kv_cache 后，首个 prefill 请求经 forward_prefill（1424）→layer_call（prefill/forward.py:637）→Block.forward:743 或 forward_prefill_fast_fp8:578 的 self.ffn（无 is_decode_forward）抛 RuntimeError；异常被 PyWrappedModel::forward 转为 std::runtime_error（PyWrappedModel.cc:1234），NormalExecutor::process 全文件无捕获，NormalEngine::step() 函数级 try（824）catch 后 895 行重抛，loop() while 体（686-692）无捕获→服务线程带异常退出且无任何重启机制（loop_thread 仅在 startLoop:666 创建、stop:674 仅 join），引擎永久停止调度。
首个 prefill 请求的守卫异常逃逸出引擎 loop()，服务线程永久终止：实例无法再服务任何请求（含 decode 与在途请求，客户端挂起直至超时），无逐请求报错兜底；监督重启后启动预热再次全绿通过，下一 prefill 请求再次杀死实例，形成崩溃循环，直至改配置重启。不产生静默错误结果；DECODE 角色实例与 auto 默认配置不受影响。
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:348：if self.strategy_name == "fp8_per_block_ep_low_latency" and not (
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:353："decode; use fp8_per_block_ep_normal for prefill"
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_per_block.py:195：or config.moe_strategy == "auto"
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:232：moe_quant_method = "FP8_PER_BLOCK"
rtp_llm/models_py/model_desc/deepseek_v4_model.py:378：args.moe_config = moe_config
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/deepep_normal_router.py:49：checker.check(not resolver.use_low_latency(config))
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/deepep_low_latency_router.py:55：checker.check(resolver.use_low_latency(config))
rtp_llm/config/server_config_setup.py:144：moe_config.use_deepep_low_latency = deep_ep_config.use_deepep_low_latency
rtp_llm/config/server_config_setup.py:230：use_deepep_low_latency = True
rtp_llm/cpp/models/PyWrappedModel.h:346：init_resources.is_decode_role = (params.parallelism_config.role_type == RoleType::DECODE);
rtp_llm/models_py/modules/dsv4/block.py:743：ffn_out = self.ffn(x_pre, ffn_in_ids) # [T, dim]
rtp_llm/models_py/modules/dsv4/block.py:578：ffn_out = self.ffn(x_pre, input_ids)
rtp_llm/models_py/modules/dsv4/block.py:479：ffn_out = self.ffn(x_pre, input_ids, is_decode_forward=True)
rtp_llm/models_py/model_desc/deepseek_v4_model.py:1424：return forward_prefill(
rtp_llm/cpp/normal_engine/NormalEngine.cc:220：if (runtime_config.warm_up && (!model_config.mm_model_config.is_multimodal)
rtp_llm/server/server_args/engine_group_args.py:25："--warm_up",
rtp_llm/server/server_args/engine_group_args.py:29：default=True,
rtp_llm/cpp/normal_engine/NormalEngine.cc:389：return prefillWarmUp(params);
rtp_llm/cpp/normal_engine/NormalEngine.cc:442：executor.reset(new NormalExecutor(params, nullptr, true, false, 0, mla_ops_type_));
rtp_llm/cpp/models/PyWrappedModel.h:341：if (params.kv_cache_layer_layout.has_value()) {
rtp_llm/models_py/model_desc/module_base.py:101：self.kv_cache = init_resource.kv_cache
rtp_llm/models_py/model_desc/deepseek_v4_model.py:1370：if self.kv_cache is None:
rtp_llm/models_py/model_desc/deepseek_v4_model.py:1381：return PyModelOutputs(hidden)
rtp_llm/models_py/model_desc/deepseek_v4_model.py:931：if getattr(moe_layer, "strategy_name", None) in (
rtp_llm/cpp/normal_engine/NormalExecutor.cc:324：model_output = std::move(model->forward(model_input));
rtp_llm/cpp/models/PyWrappedModel.cc:1234：throw std::runtime_error(std::string("pybind11 error during forward call on Python instance: ") + e.what());
rtp_llm/cpp/normal_engine/NormalEngine.cc:895：throw;
rtp_llm/cpp/normal_engine/NormalEngine.cc:687：auto status = step();
rtp_llm/cpp/normal_engine/NormalEngine.cc:666：loop_thread_ = autil::Thread::createThread(std::bind(&NormalEngine::loop, this), "normal_engine_loop");
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:27：moe_quant_method = "FP8_FP4"
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:121：use_deepep_low_latency=False,
rtp_llm/models/deepseek_v4.py:427：data_type=torch.int8,
rtp_llm/models/deepseek_v4.py:595：weights += self._build_routed_experts_stacked(layer_id)
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/layer.py:234：moe_quant_method = "FP8_FP4"
rtp_llm/models_py/modules/factory/fused_moe/strategy_registry.py:28：return quant_method != "FP8_FP4" or not isinstance(strategy, MoeStrategy)
rtp_llm/models_py/model_desc/deepseek_v4_model.py:699：self._v4_args.is_decode_role = self._is_decode_role
rtp_llm/models_py/model_desc/deepseek_v4_model.py:761：self.v4 = V4Transformer(self._v4_args, mw=self.weight)
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:203：self.is_decode_role = bool(is_decode_role)
rtp_llm/models_py/modules/factory/fused_moe/utils/fp8_fp4/chunked_layer.py:235：self.strategy_name = self.moe.strategy_name
rtp_llm/models_py/modules/factory/fused_moe/README.md:89：Use check_conditions() for strategy-specific constraints; the base
rtp_llm/server/server_args/moe_group_args.py:61：help="设置为 True 以启用 DeepEP 的低延迟模式。默认值为 None，允许自动配置。",
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/strategy/fp8_fp4.py:30：checker.check(config.moe_quant_method == "FP8_FP4")
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/routers/pure_tp_router.py:66：checker.check(resolver.use_all_gather(config))
rtp_llm/config/server_config_setup.py:109：and not deep_ep_config.use_deepep_low_latency
rtp_llm/server/server_args/scheduler_group_args.py:10："--use_batch_decode_scheduler",
rtp_llm/server/server_args/scheduler_group_args.py:14：default=False,
rtp_llm/cpp/normal_engine/NormalEngine.cc:824：absl::Status NormalEngine::step() try {
rtp_llm/cpp/normal_engine/NormalEngine.cc:388：if (pd_sep_config.role_type == RoleType::PDFUSION || pd_sep_config.role_type == RoleType::PREFILL) {
rtp_llm/cpp/normal_engine/NormalExecutor.cc:128：std::make_optional(is_propose ? cache_manager->getMTPModuleGroupedCacheLayerLayout(propose_model_index) :
rtp_llm/cpp/normal_engine/NormalExecutor.cc:130：std::nullopt,
rtp_llm/models_py/modules/dsv4/prefill/forward.py:637：h = layer_call(
建议：在 ChunkedFp8Fp4MoeLayer.init（is_decode_role 与 strategy_name 均已可得）对『LL 策略 + 非 decode 角色』直接拒绝，或 auto 模式下自动回落 fp8_per_block_ep_normal 并记录告警；亦可将 is_decode_role 注入 Fp8Fp4MoeRuntimeConfig 使 LL 策略 check_conditions 角色感知——令不兼容配置在启动期失败，而非在服务期抛出杀死引擎 loop 线程的异常。
评审版本：f8c79ba72870
dskv4-flash topp pruning