Repository navigation
[Feature] Reuse LMDeploy KPool kernels in GLM-5.3 flash indexer - #2269
qescccczmr wants to merge 7 commits into
Conversation
qescccczmr
changed the title
This comment was marked as off-topic.
This comment was marked as off-topic.
Sorry, something went wrong.
qescccczmr
changed the title
PR #2269 自查summary
Main Flowchart before this PR以下展示目标配置 graph TD
subgraph B_CFG["xtuner.v1.model.moe.glm53.nope_dsa_mla.NoPEDSAMLAConfig.build"]
BC["校验冻结约束并构造 NoPEDSAMultiLatentAttention"]
end
subgraph B_INIT["xtuner.v1.model.moe.glm53.nope_dsa_mla.NoPEDSAMultiLatentAttention.__init__"]
BI["构造 KPoolIndexer；冻结其参数"]
BA["xtuner.v1.ops.sparse_mla.get_sparse_mla"]
end
subgraph B_KINIT["xtuner.v1.model.moe.glm53.nope_dsa_mla.KPoolIndexer.__init__"]
BK["xtuner.v1.ops.sparse_mla.get_kpool_topk_indices：解析 selector callable"]
end
subgraph B_SEQ["xtuner.v1.data_proto.sequence_context.SequenceContext.split"]
BS["保留全局 cu_seq_lens；记录本 shard 起点"]
end
subgraph B_ATTN["xtuner.v1.model.moe.glm53.nope_dsa_mla.NoPEDSAMultiLatentAttention.forward"]
BF["计算 absorbed Q/KV 与 q_resid"]
BN["torch.no_grad 内调用 self.indexer"]
BM["self.sparse_mla_func 使用 int32 token IDs"]
BO["value/output projection；返回 AttnOutputs"]
end
subgraph B_KFWD["xtuner.v1.model.moe.glm53.nope_dsa_mla.KPoolIndexer.forward"]
BP["wq_b；wk + k_norm；compression gate"]
BW["weights_proj 在原 dtype 投影，结果再转 FP32"]
BT["调用 self._topk_indices_fn"]
end
subgraph B_OP["xtuner.v1.ops.sparse_mla.kpool.kpool_topk_indices"]
BG["gather_for_sequence_parallel：先 gather K/gates"]
BB["build_pools：每文档 weighted pooling，输出原 dtype"]
BR["pool_causal_ranges：全局 pool 范围"]
BQ["tilelang_indexer_topk_from_ranges：BF16 score / Top-K"]
BE["expand_pools_and_tail：排序 pool，固定列放 tail，再 pad"]
end
BC --> BI
BI --> BK
BI --> BA
BK --> BF
BS -.-> BF
BF --> BN
BN --> BP
BP --> BW
BW --> BT
BT --> BG
BG --> BB
BB --> BR
BR --> BQ
BQ --> BE
BE --> BM
BA -.-> BM
BF -.-> BM
BM --> BO
原类/权重归属、 Main Flowchart after this PR图示的是 已提交代码。黄色为新增或修改节点，红色为 A1 涉及的模型类内后端分叉；不是将建议重构画成已完成。使用 graph TD
subgraph A_CFG["xtuner.v1.model.moe.glm53.nope_dsa_mla.NoPEDSAMLAConfig.build"]
AC["冻结约束与 attention 构造保持原样"]
end
subgraph A_INIT["xtuner.v1.model.moe.glm53.nope_dsa_mla.NoPEDSAMultiLatentAttention.__init__"]
AI["构造 KPoolIndexer；冻结其参数"]
AA["xtuner.v1.ops.sparse_mla.get_sparse_mla"]
end
subgraph A_KINIT["xtuner.v1.model.moe.glm53.nope_dsa_mla.KPoolIndexer.__init__"]
AK["保留 get_kpool_topk_indices 与 tilelang 名称"]
AE["新增 _uses_lmdeploy_kpool：geometry / backend / balance 判断"]
end
subgraph A_SEQ["xtuner.v1.data_proto.sequence_context.SequenceContext.split"]
AS["保留全局 cu_seq_lens；记录本 shard 起点"]
end
subgraph A_ATTN["xtuner.v1.model.moe.glm53.nope_dsa_mla.NoPEDSAMultiLatentAttention.forward"]
AF["计算 absorbed Q/KV 与 q_resid"]
AN["torch.no_grad 内调用 self.indexer"]
AM["self.sparse_mla_func 使用 int32 token IDs"]
AO["value/output projection；返回 AttnOutputs"]
end
subgraph A_KFWD["xtuner.v1.model.moe.glm53.nope_dsa_mla.KPoolIndexer.forward"]
AP["计算 Q、projected K 与 gate 参数"]
AV{"_uses_lmdeploy_kpool 且 runtime 可用？"}
AP32["FP32 layer_norm 后 cast；fp32_linear 禁用 autocast"]
APOLD["原 k_norm 与 weights_proj 路径"]
AT["调用同一个 self._topk_indices_fn"]
end
subgraph A_OP["xtuner.v1.ops.sparse_mla.kpool.kpool_topk_indices"]
AD{"再次检查 geometry / runtime / scorer / balancing"}
AOLD["原 build_pools + TileLang + expand_pools_and_tail"]
end
subgraph A_LMD["xtuner.v1.ops.sparse_mla.lmdeploy_kpool._kpool_topk_indices"]
AG["全局 gather K/gates；build_pool_index"]
AB["_compress_storage：LMDeploy decode compression + FP8 values/scales"]
AR["pool_causal_ranges；_score_chunk_size 计算真实 padded storage 上限"]
AQ["分块 rotate_kpool_query / FP8 quantize；合并 head 权重与 scales"]
AL["DeepGEMM.fp8_mqa_logits；kpool_select_groups_cuda"]
AX["保留 group 顺序；compact tail；加文档 offset；pad int32 IDs"]
end
AC --> AI
AI --> AK
AI --> AA
AK --> AE
AE --> AF
AS -.-> AF
AF --> AN
AN --> AP
AP --> AV
AV -->|是| AP32
AV -->|否| APOLD
AP32 --> AT
APOLD --> AT
AT --> AD
AD -->|LMDeploy| AG
AD -->|fallback| AOLD
AG --> AB
AB --> AR
AR --> AQ
AQ --> AL
AL --> AX
AX --> AM
AOLD --> AM
AA -.-> AM
AF -.-> AM
AM --> AO
classDef changed fill:#fff3cd,stroke:#b58100,stroke-width:2px;
classDef review fill:#ffe8e8,stroke:#cb2431,stroke-width:2px;
class AP32,AD,AG,AB,AR,AQ,AL,AX changed;
class AE,AV review;
代码对应：模型投影位于 核心原理
核心实现与对应的单测抽象与信息隐藏评估总体：operator Module 已具有实际 Depth，但 Indexer 的数值准备与后端选择尚未形成一个完整 Seam。
A1 的解决方向：把 numerical preparation 与 selection 收敛在同一深 Module 的 Interface 下，复用一个实现选择结果；模型保留参数所有权和 checkpoint key。这样改变 geometry、fallback 或 dtype 规则时，修改与测试能集中，改善 Locality/Leverage。具体伪代码与 R1–R3 对应放在下一节；不增加无必要的抽象基类。 新增功能守则只评估本 PR 新增接触点。
A1 解决方案伪代码（仅表示职责位置，非已提交接口）： # 既有模型：保留原有参数注册位置与权重 key。
# 以一个 callable 替代旧 selector 字段，无需再保存 feature flag。
self._run_indexer = build_kpool_runner(backend, geometry, options)
def forward(hidden_states, q_resid, seq_ctx):
q = self.wq_b(q_resid).view(...)
projected_k = self.wk(hidden_states)
return self._run_indexer(
hidden_states, q, projected_k, seq_ctx,
norm=self.k_norm, head_projection=self.weights_proj,
compression_parameters=materialized_parameters,
)
# operator Module 内部：两个真实 Adapter 可用函数组合实现，
# 必要时使用纯数据 record；不要求新增类或迁移模型参数。
def run_indexer(...):
adapter = resolve_adapter(q.device) # 同一次决策控制 prepare 与 select
k, head_weights, gates = adapter.prepare(...)
return adapter.select(q, k, gates, head_weights, ...)这里的 公开接口的线性业务流程评估
有收益的下沉示意： def _kpool_topk_indices(...):
pools = _prepare_pools(...) # gather + 文档 layout + compression
groups = _score_and_select(...) # causal ranges + budget + FP8/scales + MQA
return _expand_selected(...) # group order + tail + global offset + paddingdeletion test：只有能隐藏上述整组规则、改善 Locality 的 private Module 才值得保留； 单测建议T1 — P2：新空 query 测试在 fallback 环境不成立，建议合入前修复位置： CUDA 可用但缺少 LMDeploy，或设备为 SM80 时，dispatcher 合法进入原路径；原 已在指定远端以 CPU、强制 availability=False 的最小调用核实同一 fallback 检查： 这验证的是 fallback 控制流和元数据矛盾，未声称已构建“CUDA 有但 LMDeploy 未安装”的完整测试环境。 建议：LMDeploy 专属测试的适用条件应与真实 Adapter 能力相同；公共 Interface 的 empty-input 行为则使用合法全局元数据另行定义。直接执行 MQA 的测试也检查完整能力。现有双卡测试的 guard 可作为写法参照。 # LMDeploy 专属 case；不改变生产 fallback 的语义。
if not cuda_available or not lmdeploy_kpool_available(device):
pytest.skip("Requires compatible LMDeploy/DeepGEMM and SM90+")
# 保留该 Adapter 的 empty fast-path 检查。
...
# 公共 Interface 的 fallback case：使用与全局 context 一致的 K/gates。
with force_fallback():
result = kpool_topk_indices(valid_inputs, valid_seq_ctx, ...)
assert_expected_causal_ids(result)T2 — key norm 精度的独立断言不足
# 使用非平凡 norm weight/bias；autocast on/off 分别覆盖。
z = projected_keys.double()
mu = z.mean(dim=-1, keepdim=True)
var = ((z - mu) ** 2).mean(dim=-1, keepdim=True)
expected_k = ((z - mu) / sqrt(var + eps) * norm_weight + norm_bias)
assert_bf16_close(captured_keys, expected_k.to(output_dtype))这是补足 Interface 数值契约的建议，未发现当前 key norm 公式错误；独立高精度参考应允许最终 BF16 rounding 的合理误差。 T3 — 其他覆盖范围
示意参数表： @pytest.mark.parametrize("heads,topk", [(32,2048), (64,2048), (128,8192)])
def test_supported_geometry_through_public_interface(heads, topk):
require_matching_adapter_capability(heads, topk)
assert_causal_selection_and_reference_scores(...)其他 issuesC1 — 依赖能力探测尚不完整，属于兼容性加固建议
未识别到所声明兼容 LMDeploy revision 的实际失败，不能将其写成已发生的生产回退 bug。 建议与 A1 一起在一个 Module 内解析并验证完整 operator bundle，改善 Locality；缓存 callable/module，不把可变预算值永久冻结。真正 operator 执行异常仍应抛出。 def resolve_lmdeploy_operators(device):
if not supported_device(device):
return None
try:
operators = import_all_operators_actually_used()
validate_callable(operators.mqa_layout_probe)
validate_budget_configuration_present()
return operators
except known_dependency_or_import_errors:
return None已有证据与范围，不混写成新增通过记录
收尾：本次交付为自查文档和 T1 最小复现物料，未修改被审查的源代码，未新增仓库文档目录。 合入建议仍为：修复 T1；按本次明确的零污染标准收敛 A1，或在设计审查中明确处理该不符合项。其他条目区分为覆盖建议与兼容性加固，没有发现的错误不作结论。 |
|
Updated in 4385e729.
The PR description now includes the fresh three-step MoE-only comparison on the same 1650 / 8×H200 host:
BF16/BF16 had 8.5233% lower mean These are three-step, independently generated trajectories on a five-layer reduced checkpoint, comparing complete KPool pipelines rather than quantization alone. The indexer stayed frozen and rewards/losses/gradient norms were zero, so this is numerical mismatch evidence, not a learning or model-quality claim. The implementation AST matches the measured precision-switch experiment apart from docstrings; no new GRPO run on this final commit is claimed. |
qescccczmr
changed the title
|
上面均在本地复现，且审核代码没有发现问题。
|
summaryNo blocking findings remain in the KPool follow-up. The BF16 startup option now selects serving compression and FP32 model preparation on supported SM90+ devices; the BF16 TileLang scorer remains. LMDeploy's established arithmetic is unchanged; tests pin it down. Main Flowchart before this PRflowchart TD
subgraph XTuner_KPoolIndexer_forward
A[BF16 override] --> B[BF16 head projection then FP32 cast]
end
subgraph XTuner_kpool_build_pools
B --> C[FP32 softmax then BF16 probabilities/products]
end
subgraph XTuner_kpool_topk_indices
C --> D[TileLang score / causal Top-K / token expansion]
end
Main Flowchart after this PRflowchart TD
subgraph XTuner_KPoolIndexer_forward
A[BF16 override + supported serving operators] --> B[FP32 normalization and fp32_linear]
end
subgraph XTuner_kpool_build_pools
B --> C[LMDeploy decode compression / FP32 intermediates / BF16 store]
end
subgraph XTuner_kpool_topk_indices
C --> D[TileLang score / causal Top-K / token expansion]
end
style B fill:#d9f2d9,stroke:#267326
style C fill:#d9f2d9,stroke:#267326
核心原理BF16 storage does not imply BF16 pooling products or head projection. For matching inputs, keep those computations in FP32 until pool storage. Indexer IDs remain frozen and nondifferentiable; no new compression backward is introduced. 核心实现与对应的单测
On 10.102.215.14 / H200: XTuner 62 passed / 6 skipped (BF16 run), SP 2 passed, LMDeploy precision 29 passed. Compact MQA and FP8-only public paths account for skips. Fixed real-model replay has 32632 predefined valid scores, zero nonfinite values, identical pool/head values and A/A output, score relative L2 2.2019015228460665e-7, maximum absolute error 3.814697265625e-5, and 0/32 different Top-K sets. This is an operator check, not end-to-end KL acceptance. 抽象与信息隐藏评估The existing adapter module earns its depth: compression, optional dependencies and storage dtype stay at the existing seam. No new production class/module is added. The interface adds a dtype to an existing private helper and a precision-specific availability query, providing locality without moving serving scheduler state into training. 新增功能守则The existing model has one preparation condition changed; no feature field or numerical implementation is added to it. FP32 projection uses the existing 公开接口的线性业务流程评估Projection -> pool construction -> scoring -> causal selection -> token expansion remains linear. Fallbacks and public backend names remain available. FP8 compression calls retain their old call form for compatibility with older LMDeploy versions. 单测建议Keep the new real-shape and boundary regressions. The paired replay is the acceptance surface for the local numerical target; global KL, multi-update training, cache lifecycle and throughput remain separate checks. 其他 issuesNo additional blocking finding. The remaining FP32 scorer reduction-order error is explicitly reported; this change does not claim bitwise score equality or whole-model KL <= 2e-4. Non-SM90 / optional-dependency fallback paths retain legacy arithmetic and are outside this parity claim. |
Root Cause
XTuner's GLM-5.3 flash KPool indexer used BF16 pool scoring while LMDeploy uses rotated FP8 pools/queries and DeepGEMM MQA scoring. Different precision boundaries and selected-group/tail ordering can change the sparse attention inputs and contribute to rollout/training mismatch.
Fix
Reuse LMDeploy's decode compression, query rotation/FP8 quantization, MQA scoring, group selection and token expansion for supported GLM-5.3 geometry: dimension 128, 32/64/128 heads, pool size 4, Top-K 2048/8192.
Preserve
indexer_backend="tilelang"and existing function names. Check GPU/API availability at forward time and retain the original implementation/projection arithmetic when optional LMDeploy/DeepGEMM dependencies or supported hardware are unavailable. Alternative geometries/scorers/selectors and SP query balancing retain their previous path.Match FP32 key normalization and head-gate projection, explicitly disabling autocast for the latter via
fp32_linear.Gather keys/gates before pooling under SP, retain local query ownership/global token IDs, and preserve serving group order and compact causal tails.
Chunk query rotation/quantization and bound logits storage using the active MQA API's actual layout.
Keep this adapter forward-only for the frozen DSA indexer. Remove the unused differentiable compression API, custom STE backward and gradient-only tests. The actual
_compress_storageand_kpool_topk_indicespath remains in place.Keep operator contracts in the existing KPool docstrings and experimental evidence in this PR description. The final diff consists of three implementation files and two operator tests, using the existing source/test directories.
KPool precision switch
Set
XTUNER_GLM53_KPOOL_BF16before training workers start, including propagation through the Ray runtime environment when applicable:01Cached precision-specific availability checks keep model preparation and pooling consistent. The BF16 scorer remains TileLang; the FP8 path retains the serving adapter. Keep
indexer_backend="tilelang"; no backend renaming or extra model configuration is needed. Restart workers to change the value. MoE-only FP8 is unaffected.This is an XTuner-only switch. The BF16/BF16 experiment below additionally used a separate LMDeploy BF16 cache/scorer patch with
LMDEPLOY_GLM53_KPOOL_BF16=1; that inference-side patch is not part of this XTuner PR. Setting only the XTuner variable does not make a default FP8 rollout engine BF16.BF16 pooling/head precision follow-up
At
354a1fd6, the BF16 override reuses LMDeploy decode compression rather than rounding probabilities and products to BF16 before pooling. It also promotes head weights before the projection through the existing autocast-safefp32_linear. No backend or environment variable was renamed. The frozen indexer still has no compression backward.Paired with LMDeploy #5115 (
1c52dae9), a fresh same-input replay on 10.102.215.14 / H200 reproduced:Inputs came from a captured 4,096-token real-checkpoint prefix; both arms consume identical tensors. Causal slots were fixed from sequence metadata before comparison, excluding 136 expected future
-infentries rather than intersecting finite outputs. The residual is FP32 scorer reduction order, not a claim of bitwise scores or end-to-end KL at e-7.Fresh follow-up checks: XTuner 62 passed / 6 skipped with the BF16 option, two-GPU SP 2 passed across BF16/FP8, and LMDeploy precision tests 29 passed. Skips cover compact MQA unavailable in the installed wheel and FP8-route-only cases under the BF16 override. Changed-file Ruff and syntax checks pass. Default FP8 compression keeps its existing call form.
Usage: install both PRs into their corresponding environments, preserve
indexer_backend="tilelang"andfreeze_dsa_indexer=True, and export both startup variables before launching/restarting all training and rollout workers:The complete MoE-only recipe additionally uses
FP8=1 ROLLOUT_FP8=1 LMDEPLOY_FP8_MOE_ONLY=1; our current diagnostic recipe also keepsLMDEPLOY_MLA_DISABLE_DENSE_PREFILL=1. Propagate these through Ray runtime environments. The existing SwiGLU/mHC patches are separate from this PR. A whole-modelmismatch/mismatch_k3_kl <= 2e-4remains an end-to-end validation target, not a result established by this operator replay.Review: legacy MQA allocation budget
Addresses the legacy-budget review: compact columns use
align(N, 256); legacy absolute columns usealign(N + 256, 64), including the guard columns. Both APIs pad query rows to128 / heads. Chunking now accounts for the selected layout rather than applying the compact formula to both.The regression checks
logits.untyped_storage().nbytes(), including short unaligned chunks. For 4096 tokens / 1024 pools / 32 heads / a 1-MiB budget, the old 256-row chunk allocated 1,310,720 bytes. The corrected 204-row chunk allocates 1,044,480 bytes, within 1,048,576 bytes.Scope
This reuses serving operators; it does not claim complete prefill/decode scheduling or full-model log-probability parity. Training compresses closed pools with
mode="decode", while serving initially builds prompt pools with onlineextendcompression.SequenceContextlacks the prompt/chunk history needed to reproduce that schedule. The adapter uses internal LMDeploy helpers verified at3228395546dfb2f6427faf39ec0c0c285f4c378a.Top-K IDs remain nondifferentiable, the indexer stays frozen, and
freeze_dsa_indexer=Falseremains unsupported. This PR adds no compression backward or end-to-end indexer training. Existing attention/MoE backward is unchanged. It contains no mHC, KDA QKV, SiLU or MoE W2 modifications.Test Plan
Precision-switch follow-up: fresh CPU/GPU regressions on the requested H200 host: 64 passed / 2 skipped with FP8 enabled, and 59 passed / 7 skipped with the BF16 override. The latter skips five FP8-route-only cases plus two unavailable compact-MQA cases. Tests verify that BF16 bypasses optional CUDA/dependency probes and preserves legacy model preparation plus selector dispatch, with CPU/CUDA autocast on/off.
Public-selector tests that specifically require LMDeploy FP8 now check the enabled/available path, including the empty-local-query fixture. This avoids running that fixture through an incompatible fallback. Changed-file pre-commit and Ruff pass; mypy checks all 431 source files.
Layout/docstring follow-up at
ffe21b9d: executable AST unchanged relative to79d33b3d; the renamedtests/ops/test_lmdeploy_kpool_sequence_parallel.pypasses a fresh two-GPU check. That revision collected 61 cases under the final filenames; the precision-switch follow-up now collects 66.H200 CPU/GPU regression at
79d33b3d: 59 passed, 2 skipped, including two-rank NCCL SP. The skipped cases require compact MQA, which the installed legacy wheel does not expose; CPU layout tests cover both contracts and real-storage GPU tests cover the available legacy API.Pool FP8 values/scales match the serving operator bitwise; independent PyTorch forward comparisons cover both compression modes and scale settings. Empty groups, architecture fallback, CPU/CUDA autocast, dependency fallback, packed causality and group/tail order remain covered.
Query FP8 payloads and folded scale weights are checked at the actual MQA call. Valid causal logits are checked against independent Torch FP8 scoring using DeepGEMM's simulated-FP8 normalized-error criterion (
5e-6).Fixed inputs, seed 123, packed lengths
(8195, 8189), heads32/dim128/Top-K2048: zero changed token IDs versus75aeaec1at row caps 1024/257 and logits budgets 512/1 MiB.Changed-file pre-commit, explicit Ruff and mypy across 431 source files pass. Full-repository pre-commit also ran in an isolated checkout; seven untouched
autotest/files have existing formatting/import issues, excluded from this change.Recorded MoE-only GRPO check
The following three-step A/B results were measured at
75aeaec1, before removing the unused backward, against pre-adapter baselinec93eb15d. Both arms ran on the same 8×H200 host with the same five-layer checkpoint, seed 123, EP4/SP1 training, TP1/EP4 rollout and 64 responses per step. Attention/dense GEMMs remained BF16, grouped MoE GEMMs used tilewise FP8, and serving's effective quantization scope wasmoe_only.Mean
mismatch/mismatch_k3_klwas 7.95% lower. Current forward parity against that measured commit is checked above. These independently generated trajectories are not a fixed-token causal ablation; rewards/losses/gradient norms were zero, so this is numerical smoke-test evidence, not evidence of learning or model-quality gains.Recorded BF16/BF16 versus FP8/FP8 MoE-only GRPO check
A separate fresh two-arm check on 2026-11-09 used the same 1650 / 8×H200 host, five-layer checkpoint, seed 123, EP4/SP1 training, TP1/EP4 rollout, 64 responses per step and length limits 2048/14336/16384. The DSA indexer stayed frozen; grouped MoE GEMMs remained tilewise FP8, ordinary attention/dense GEMMs remained BF16, and serving's effective quantization scope was
moe_only. No mHC, KDA QKV or SiLU/W2 patches were enabled.Both arms used
ffe21b9dplus this startup gate and the same LMDeploy32283955source with the MoE-only correction and the separate BF16-selection patch. The three implementation files in this follow-up have the same executable AST as that measured experiment after removing docstrings; this is recorded evidence rather than a claim of another GRPO run on the final commit.For BF16/BF16 both
XTUNER_GLM53_KPOOL_BF16andLMDEPLOY_GLM53_KPOOL_BF16were1; for FP8/FP8 both were0. Each arm completed all three steps with exit code 0 and 192 responses. The actual environments of eight training and eight inference workers per arm were checked. Serialized training configurations and source revisions/worktrees matched; recorded settings differed only in the two precision variables and output directory.BF16/BF16 had 8.5233% lower mean
mismatch/mismatch_k3_klin this run, while both remained at the e-4 level. This compares complete KPool pipelines: projection precision, pooling accumulation and group/tail ordering also differ, so the result cannot be attributed solely to FP8 quantization. The trajectories were independently generated, not paired fixed-token inputs. Rewards, losses and gradient norms were zero; three steps on a reduced model do not establish learning or full-model quality gains. The default remains FP8 to match an unmodified FP8 rollout engine; the switch preserves the BF16 option for deployments that configure both sides accordingly.Upstream CI status is tracked by the PR checks.
Five-step combined numerical check (2026-11-11, H200)
A fresh five-step check completed on 10.102.215.14 / 8 x H200, exit code 0, using the five-layer reduced GLM-5.3 Flash 25B checkpoint, GSM8K, EP4/SP1 training and TP1/EP4 rollout, prefix caching and MTP with
num_speculative_tokens=5.This is a combined recipe: BF16 KPool pooling/head alignment, FP8 MoE activation rounding, opt-in FP32 KDA chunks, and the separate mHC-post/SwiGLU alignment patches. The eight startup variables
FP8,ROLLOUT_FP8,LMDEPLOY_FP8_MOE_ONLY,XTUNER_GLM53_KPOOL_BF16,LMDEPLOY_GLM53_KPOOL_BF16,XTUNER_GLM53_LMDEPLOY_SWIGLU,XTUNER_GLM53_KDA_FP32, andLMDEPLOY_MLA_DISABLE_DENSE_PREFILLwere all1. The actual environments of 8 training workers, 8 inference-compute workers and their engine/controller wrappers were checked; the prefix-cache replay probe also passed.mismatch/mismatch_k3_klThe end-to-end <=2e-4 target was not reached (all five values remain above it). TensorBoard scalar values were cross-checked with the worker's full-precision log values. Local KPool score relative L2 at 2.2e-7 is a different metric and does not imply whole-model KL at that level.
Rewards, losses and gradient norms were zero; this is a numerical GRPO smoke test, not evidence of useful learning or nonzero parameter updates. No per-patch causal improvement is inferred from this combined run. Remaining differences, including the actual DeepEP dispatch/combine path and KDA recurrence, are still under diagnosis.