iframe-proxy

guluguluhhhh · 2026-10-19T14:51:38Z

Summary

Reduce DSv4.1 Flash decode overhead through FULL attention overlap, Reindex projection and sparse-selection fusion, and batched MTP/DSpARK preparation. Target branch: feature/dsv41-dev-main-202161013.

In the B64/rank, ctx32K, 4DP4EP workload below, the latest measured revision 667f2a13d improves full-model fixed-batch decode TPS by 5.37% and reduces MTP step time by 5.09% against the retained same-day target baseline dd6dd239a. The baseline was not rerun in this attempt. A separate, earlier timeline analysis at 438fbb2a0 shows a 10.54% reduction in the layer-count-weighted Attention interval. Accuracy results are reported separately with their own reference revision.

Attention Optimizations

FULL1/FULL2 compressor: remove the fused Global producer's 192-token cap while retaining device, layout and state checks. This enables the existing compression/norm/RoPE/cache-store fusion at B64 with MTP gamma=5, where target verification processes 384 tokens per rank.
FULL metadata and scheduling: fuse visible-length, request-ID and page-table preparation; share compatible metadata within a forward. Run QKV, Global KV production, and indexer weights/metadata on three streams with explicit joins and buffer lifetime management. Overlap FULL1 candidate selection with token selection/attention, and support Q-norm/quantization -> merged QB programmatic dependent launch (PDL).
Reindex front projection: combine FP8 Q-A/KV projection and BF16 index-head weight GEMM in one native mixed-precision kernel on supported B300 shapes. Head weights remain BF16; each head tile completes its K reduction without separate split-K scratch/reduction. DeepGEMM device headers are integrated at build time; 9b183007 adds compatibility with the locked device-header layout.
Reindex QB epilogue: fuse Main + Indexer QB GEMM with Indexer-Q RoPE, FP4 quantization and head-weight preparation using the guarded BN128 Indexer-first path. Preserve shared resident projection weights, reload handling and fallback consumers. Front/QB fusions are restricted to Reindex; FULL retains its shared-input, multistream and PDL path. Reindex scoring, TopK, index mapping and attention still execute, and the upstream O-A implementation is retained.
Sparse token selection: retain native BF16 scorer output through DeepSelect without materializing FP32 logits. Complete engine plans select directly from disposable scorer output using compact scan lengths; compatibility plans repair filler/zero-page slots in place, while borrowed logits retain the copy-based cleanup path with 2048-element tiles. Preserve finite filtering, logical-index restoration and TopKv3 fallback.
Sparse Plan preparation: reuse causal lengths across Reindex consumers, removing redundant per-layer position-add/floor-div work. Publish sorted candidate prefixes from eligible FULL producers and consume them without re-sorting or flag compaction. Other inputs retain flag compaction and bounded int16/int32 sorting fallbacks. Propagate replay-updated request validity from the CUDA Graph runner so padded requests have zero compact scan lengths without changing padded MTP/KV-write geometry. The official DeepGEMM Sparse MQA metadata operator is unchanged.
Plan overlap: prepare the reusable Plan on a separate stream after FULL candidate publication, waiting on the actual auxiliary producer and joining before Reindex consumption. Preserve forward-local lifetime, interrupted-forward cleanup and cold-capture fallback. Launch begins after candidate publication; the deferred attention-end mode has been removed.
Optional SM103 paths: add custom merged QB/KV production, TopK cleanup plus logical-to-physical slot conversion, and REUSE Q/KV/index production. These paths are disabled by default and were not enabled in the results below. REUSE attention uses the existing FlashMLA wheel. The source-pinned native FlashMLA consumer and outgoing ready-QKV PDL handshake were removed in 2d454c4f; the incoming QB PDL path is retained.

Other Decode Optimizations

MTP bookkeeping: batch stochastic/forced-accept masks into one pinned transfer and reuse complete GPU page-table snapshots, retaining mixed/partial-batch fallbacks.
DSpARK preparation: fuse proposal token IDs, positions and active-mask preparation, plus commit request/position row mapping. Preserve integer casting, duplicate-start, padding and fallback behavior. Originally added in bdb284c7 and replayed as 6df05f77f, this is included in the latest end-to-end measurement at 667f2a13d; the earlier Attention and accuracy measurements do not include it.

Configuration

Shape/device eligibility checks remain in place. Unsupported inputs and cold captures retain the corresponding fallback paths. The FULL model path does not enable DeepGEMM's process-global PDL policy itself.

End-to-End Decode Performance

Four B300/SM103 GPUs, TP1/DP4/EP4, B64 per rank (256 concurrent requests), ctx32768, DSpARK gamma=5, Sparse MQA and CUDA Graph enabled. Each request produces 192 tokens in 32 MTP rounds, with six output tokens per round. Target verification processes 384 tokens/rank and draft proposal processes 320 tokens/rank. Both arms use asynchronous engine execution and FAKE_BALANCE_EXPERT=0.

This measures the complete model's decode execution through the RTP engine, including draft proposal, target verification, draft commit, sampling and engine overhead. The fixed-batch workload uses synthetic KV, skips real prefill and forces full draft acceptance. It is not an online-serving throughput measurement including prefill, queueing or network latency, and does not measure a real draft acceptance rate.

The latest run on 2026-10-21 measures PR 667f2a13d using FULL multistream + PDL. Target dd6dd239a and previous PR 438fbb2a0 below are retained same-day measurements, not reruns in this attempt. All three use internal revision dabfb8112, the same physical GPUs, workload and DeepGEMM runtime. The PR builds include the locally approved BUILD-only repair for nested select(); inference logic is unchanged by that repair. The three native library hashes match between the two PR runs. These results cover the listed revisions, not subsequent PR or target-branch changes.

Configuration	Per-request TPS	Aggregate TPS	MTP Step (ms)	TPS Gain
Target baseline `dd6dd239a`, earlier same day	264.79	67,786.63	22.659	Reference
Previous PR `438fbb2a0` + BUILD repair, earlier same day	276.22	70,711.73	21.722	+4.32%
Latest measured PR `667f2a13d` + BUILD repair, multistream + PDL	279.00	71,424.27	21.505	+5.37%

At 667f2a13d, MTP step time decreases by 1.154 ms (5.09%) versus the retained baseline. Compared with the previous PR run, TPS increases 1.01% and step time decreases 0.217 ms. The latest three unprofiled rounds measure 279.0010, 278.9747 and 279.0483 per-request TPS (range 278.97-279.05). The runs occurred at different times, so the incremental gain is not an isolated on/off measurement of the DSpARK fusion.

Each run uses two warmup rounds, three unprofiled measurement rounds and one separate profiling round. Exclude the first six-token round from each request's engine timing: token_ms = (cost_time_us - first_token_cost_time_us) / 1000 / 186. Take each rank's request mean, then the slowest rank, then the median across the three measurement rounds. Per-request TPS is 1000 / token_ms, aggregate TPS is 256 * per-request TPS, and MTP step time is 6 * token_ms. Profiler timings do not enter these results.

The PR run enables FULL multistream, its Q-norm-to-QB PDL trigger and the DeepGEMM PDL policy; baseline has these disabled. Other features retain their branch defaults, including disabled custom QB/KV/slots/ready-QKV paths. The measured gain is the combined result of the active changes and PDL policy, not an isolated attribution to one operator.

Attention Performance

The following historical timeline compares target dd6dd239a with PR 438fbb2a0, using the same model and workload as the end-to-end comparison but a separate profiling round. It is not a measurement of the later 667f2a13d revision.

The table is rank 0's mean per-layer GPU Attention interval over the first three captured target-verify iterations, grouped by Attention type. The interval starts after input HC/norm and ends at Attention output projection completion, before the FFN mHC transition. It excludes MoE and is not an end-to-end MTP step measurement.

Attention Type	Layers	Baseline (us)	Optimized (us)	Time Reduction
SWA	2	97.86	99.29	-1.45%
FULL2	3	449.17	292.24	34.94%
FULL1	1	477.58	360.80	24.45%
Reindex	4	229.85	187.66	18.35%
Reuse	30	126.61	128.06	-1.14%
Layer-count-weighted mean	40	168.47	150.72	10.54%

Positive reduction means faster. The improvements are concentrated in FULL and Reindex; SWA and Reuse show small slowdowns in this capture. The optional custom QB/KV/slots/ready-QKV paths were disabled, so these measurements do not establish their benefit.

Recorded Operator Measurements

Historical B300 microbenchmarks from the Reindex/Plan changes, recorded in the integration comment at 801eb19c. These are separate from the layer timeline above and are not a fresh full-PR end-to-end measurement.

Scope	Shape / Input	Before (us)	After (us)
Sparse `select_tokens`	M384, 2048 candidate blocks, random scores	44.06	26.16
Sparse `select_tokens`	Same shape, near-equal scores	41.61	25.94
Plan preparation	M384, K2048, domain4224	27.21	12.17
Full Plan, including unchanged metadata operator	Same shape	44.88	30.13

The selection benchmark alternates CUDA Graph variants with eight rotating inputs. Operator timings and overlap savings are not added together or treated as request-throughput gains.

Accuracy Evaluation

The completed evaluation covers 262/262 requests, with exit code 0, no failed requests and no surviving test processes. It compares the source branch's default path at 801eb19c -> 438fbb2a under the same evaluation setup. This reference differs from the target-branch performance baseline above.

Metric	Previous Revision `801eb19c`	Tested Revision `438fbb2a`	Change
GPQA, Markdown-compatible scoring	75.76%	74.75%	-1.01 percentage points
MRCR 16K-32K	65.56	65.56	Unchanged
MRCR 64K-128K	68.86	68.86	Unchanged

All 64 MRCR raw outputs are identical between revisions. GPQA has a net decrease of two correct answers: five changed from correct to incorrect and three from incorrect to correct; truncated responses increased from four to five. These results establish unchanged MRCR outputs and a small GPQA score difference on this sample, not a blanket guarantee of zero regression.

The evaluation covers the default source-branch path; default-off kernels are not counted as validated. The accuracy table and historical Attention table do not include the later DSpARK preparation change bdb284c7. The latest end-to-end row at 667f2a13d does include that change; no new accuracy evaluation is claimed for this revision or subsequent heads.

Reproduction and Validation

The PR includes regression coverage for cache slots and metadata, fallback dispatch, projection reload, changed-input CUDA Graph replay, cross-stream dependencies, compact candidate prefixes, replay-updated request validity, and MTP masks/page tables.

The Reindex/Plan integration record reports the following validation runs through 801eb19c:

Reindex integration: 52 tests passed, with no failures, errors or skips. Coverage includes FULL/Reindex dispatch, numerical references, shared projection storage, weight reload, CUDA Graph replay, sparse masking/selection and fallback paths. The M384 projection-chain test reloads weights three times and replays the original graph.
Plan/FULL-producer integration: 54 tests passed, with no failures, errors or skips. Coverage includes the actual FULL auxiliary producer, exact serial-vs-overlap Plan comparison, four Reindex consumers, changed positions/candidates/page tables/Q/weights during replay, interrupted-forward cleanup and cold capture.
Default-on Plan overlap: 7 candidate-launch and 7 attention-end tests passed, with no failures or skips, including unset defaults and explicit opt-out.

These are historical, overlapping runs, not an additive test count. The attention-end launch mode covered in those runs has since been removed. They used source-verified existing native operator artifacts and do not establish a fresh complete native build, full checkpoint/service reload acceptance, or coverage of the later C++ MTP gatherer changes. The later model accuracy evaluation is reported separately above.

Historical FULL scheduling and PDL reproduction
Historical ready-QKV and FlashMLA integration: records an earlier implementation; its native FlashMLA consumer is no longer part of this PR.
Historical MTP preparation validation: separate kernel/contract validation. Full-model decode performance including this change is reported at 667f2a13d above; accuracy coverage remains at 438fbb2a0.

rtp-llm-review-bot · 2026-10-19T15:02:08Z

自动审查进度

项目	当前状态
版本	`9b183007e7c2`
审查	检查代码改动（本阶段 9/69）
预计还有多久	约 765–1150 分钟（同规模历史耗时估算，非承诺）
自动批准	已取得人工 CODEOWNER 授权；当前版本审查通过、问题及回复处理完成后可自动批准。

更新时间：2026-10-22 11:54 北京时间。运行期间约每 5 分钟更新；阶段任务数会随核验展开，不能当作整轮完成百分比。

rtp-llm-review-bot · 2026-10-20T00:35:52Z

+        # Producer only: QA completion still orders this kernel. The following
+        # QB may initialize early but must GDC-wait before reading Q/scales.
+        tl.inline_asm_elementwise(
+            "griddepcontrol.launch_dependents; // $0",


[P2] QB 入口级 PDL 触发先于全部 Q/SF 写入，读序安全单点依赖快照外 DeepGEMM 消费端读前 gdc.wait，门控仅检查 get_pdl() 且无构建绑定

部署方已使进程内 deep_gemm.get_pdl() 为真（生产代码从不调用 set_pdl，需运维预置），再开 DSV41_FULL_DECODE_MULTISTREAM=1 与 DSV41_FULL_DECODE_QB_PDL=1，且满足 FULL 门控（TP1、SM10x、H32/D128/TopK512、merged_qb 存在、eager 预热后，attention_v41.py:3170 进入 try_full_decode）时，FULL 源层每步解码与每次图重放都经过该 programmatic 边；若所装 DeepGEMM 构建的 PSS 消费内核在 gdc.wait 之前读取 A/scales，即持续读到未写入数据。fallback 路径（attention_v41.py:3186）不传 qb_pdl，默认两开关均 "0" 时行为与 base 一致。
base 无触发指令、无 full_overlap.py、compute_qkv 不传 pdl（base:98）、try_project_quantized_qkv 无 pdl 参数（base:85）；head 新增入口级触发（:29-39）并经 full_overlap.py:172 → compute_qkv.py:99 → query_norm_quant(pdl) → PDL=pdl（:93）接线至 merged.forward_quantized（compute_qkv.py:113）的 deep_gemm 调用，首次把消费端可执行窗口从『生产者完成/数据已写出』提前到『生产者入口、Q/SF 完全未写入』。
PDL=True 时 _query_norm_quant_kernel 在入口、任何 Y/Q/SF store（:47/:57/:65）之前发射 griddepcontrol.launch_dependents（:33），允许同流以 PSS 属性启动的 merged QB deep_gemm.fp8_fp4_gemm_nt 提前执行其 wait 之前的代码；触发与 QB GEMM 之间（compute_qkv.py:100-112）仅主机端切片/getattr/profiler 上下文与 torch.empty，无其他设备内核发射。消费端调用（_v41_merged_qb.py:65-70）不传任何同步参数，是否读前等待完全由外部库内核决定。_qb_pdl_enabled 只检查两个环境变量、merged_qb 存在与 hasattr(get_pdl) 且为真（full_overlap.py:23-33），不校验消费内核 wait 位置、版本或构建。
若决定性未知被否定（消费端读前不 wait）：QB 读到未写出的 Q/SF，q/index_q 及其后全部 attention/解码输出静默数值错误，随捕获的 CUDA Graph 持续重放，生产路径无运行时校验可拦截；若消费端如文档声称在读前 wait，则仅重叠初始化，无正确性影响。严重度按未知成立时的后果评估。
rtp_llm/models_py/modules/dsv41/_v41_query_quant.py:33："griddepcontrol.launch_dependents; // $0",
rtp_llm/models_py/modules/dsv41/_v41_query_quant.py:47：tl.store(Y + row * D + cols, y, cols < D)
rtp_llm/models_py/modules/dsv41/_v41_query_quant.py:57：tl.store(Q + row * D + cols, tl.reshape(quant, (BLOCK,)), cols < D)
rtp_llm/models_py/modules/dsv41/_v41_query_quant.py:27：PDL: tl.constexpr = False,
rtp_llm/models_py/modules/dsv41/_v41_query_quant.py:93：PDL=pdl,
rtp_llm/models_py/modules/dsv41/_v41_query_quant.py:31：# QB may initialize early but must GDC-wait before reading Q/scales.
rtp_llm/models_py/modules/dsv41/_v41_query_quant.py:26：BLOCK: tl.constexpr,
rtp_llm/models_py/modules/dsv41/_v41_query_quant.py:85：def try_project_quantized_qkv(attn, x):
rtp_llm/models_py/modules/dsv41/fp8/decode/full_overlap.py:25：os.environ.get("DSV41_FULL_DECODE_QB_PDL", "0") != "1"
rtp_llm/models_py/modules/dsv41/fp8/decode/full_overlap.py:33：return hasattr(deep_gemm, "get_pdl") and deep_gemm.get_pdl()
rtp_llm/models_py/modules/dsv41/fp8/decode/full_overlap.py:172：qb_pdl=_qb_pdl_enabled(mega),
rtp_llm/models_py/modules/dsv41/fp8/decode/compute_qkv.py:99：quantized_qkv = try_project_quantized_qkv(attn, x, pdl=qb_pdl)
rtp_llm/models_py/modules/dsv41/fp8/decode/compute_qkv.py:113：q, index_q = merged.forward_quantized(*qr_quantized)
rtp_llm/models_py/modules/dsv41/fp8/decode/compute_qkv.py:98：quantized_qkv = try_project_quantized_qkv(attn, x)
rtp_llm/models_py/modules/dsv41/fp8/_v41_merged_qb.py:65：deep_gemm.fp8_fp4_gemm_nt(
rtp_llm/models_py/modules/dsv41/fp8/attention_v41.py:3170：output = try_full_decode(
docs/benchmark/dsv41_full_layer_optimize.md:111：DeepGEMM QB consumer must launch with programmatic serialization and wait before
docs/benchmark/dsv41_full_layer_optimize.md:371：cudaGridDependencySynchronize(), while actual A/B TMA loading follows it.
docs/benchmark/dsv41_full_layer_optimize.md:129：2.8.2+30e85d9.cu132, and a FlashMLA wheel exposing Mega Attention.
deps/requirements_torch_gpu_cuda13.txt:40：http://artlab.alibaba-inc.com/1/pypi/rtp_llm/deep_gemm/deep_gemm-2.8.2%2Bad62a90.cu132-cp310-cp310-linux_x86_64.whl
rtp_llm/test/perf_test/dsv41_full_decode_pdl.py:191：raise RuntimeError(f"Expected one Q norm -> QB edge, got {qb_edges}")
rtp_llm/test/perf_test/dsv41_full_decode_pdl.py:268：torch.testing.assert_close(value, ref, rtol=0, atol=0, msg=key)
rtp_llm/test/perf_test/dsv41_full_decode_pdl.py:105：deep_gemm.set_pdl(library_pdl)
rtp_llm/models_py/modules/dsv41/fp8/test/test_v41_global_overlap.py:572：if not hasattr(deep_gemm, "get_pdl") or not deep_gemm.get_pdl():
rtp_llm/models_py/modules/dsv41/fp8/test/test_v41_global_overlap.py:192：dg.set_pdl.assert_not_called()
rtp_llm/models_py/bindings/cuda/kernels/scaled_fp8_quant.cu:106：asm volatile("griddepcontrol.launch_dependents;");
rtp_llm/models_py/bindings/cuda/kernels/scaled_fp8_quant.cu:83：asm volatile("griddepcontrol.wait;");
rtp_llm/models_py/bindings/cuda/launch_utils.cc:14：enablePDL = !StaticConfig::user_disable_pdl;
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_hybrid_executor.py:853：down_input_fp8, down_input_scale = trt_fp8_quantize_128(down_input, False)
rtp_llm/models_py/modules/factory/fused_moe/impl/cuda/executors/deepgemm_hybrid_executor.py:864：m_grouped_fp8_gemm_nt_contiguous(
rtp_llm/models_py/modules/dsv41/utils.py:84：deep_gemm.fp8_fp4_gemm_nt(
rtp_llm/models_py/modules/dsv41/_v41_query_quant.py:105：projected = linear(x)
docs/benchmark/dsv41_full_layer_optimize.md:44：| deep_gemm.get_pdl() | Must already be true for that producer trigger to be enabled |
待确认：部署所用 DeepGEMM 构建的 fp8_fp4_gemm_nt 在 get_pdl()=true 时是否以 programmatic 序列化属性启动、及其 gdc.wait 相对首次 A/scales 读取的位置——库源/二进制不在快照（3rdparty/deep_gemm 仅两个无关补丁），仓钉 wheel ad62a90（deps:40）与文档验证构建 30e85d9（docs:129）哈希不同，门控（full_overlap.py:33）无任何版本绑定；这是唯一能坐实或推翻本失败的事实

建议：在 _qb_pdl_enabled 中把门控从 hasattr+get_pdl() 收紧为绑定经核验的 DeepGEMM 版本/构建白名单（与 runner 验证产物一致），或启用前要求复现 runner 式图边＋PTX 审计与逐位探针通过；或将触发点后移至 Q/SF store 之后（对齐 scaled_fp8_quant.cu:106 的仓内 post-store 纪律），使消费端即使缺读前 wait 也退化为可见性竞态而非读取未写入数据；维持双开关默认关闭。

评审版本：597c930afcb4

rtp-llm-review-bot · 2026-10-20T00:35:55Z

+    with torch.cuda.device(x.device):
+        capturing = torch.cuda.is_current_stream_capturing()
+    streams = (_GLOBAL_COMPUTE_STREAMS, _INDEXER_STREAMS)
+    if capturing and any(device not in cache for cache in streams):


[P2] 新增 FullOverlapCPUTest.test_cold_capture_and_failed_producer 在无可见 GPU 环境因 try_full_decode 对 Mock 缓存流重建真实 CUDA 流而确定性 ERROR

在无可见 GPU 的机器（纯 CPU 节点或隐藏 CUDA 设备、torch 可导入）以 pytest/unittest 直接运行 head 的 test_v41_global_overlap.py（支持 main 直接执行，640 行）；即作者文档 268-275 行自述采用的 GPUs hidden 验证实践。bazel 目标因 exec_properties 绑定 GPU 节点（BUILD:97-100）且 pytest_main 要求足量 CUDA 设备（pytest_main.py:45-57），默认 bazel 路径不进入 CPU-only 场景。
base 的 decode/ 目录无 full_overlap.py，base 测试文件仅含 GlobalOverlapCPUTest（31 行）与 GlobalOverlapCudaTest（177 行）两类；head 新增该模块、FullOverlapCPUTest 与该用例，并把侧流获取从姊妹函数 start_global_decode 的『缓存命中即直取』（head global_overlap.py:87 保持不变）改为『无条件经 _get_process_cuda_stream 复核并可重建真实流』。attention.py 的 _get_process_cuda_stream 本身 base→head 未变（两侧 165-175 行逐行一致），变化在于新调用方使其对 Mock 缓存命中也生效。
head 新增 try_full_decode 的冷捕获检查（full_overlap.py:117）仅在 device 缺失于任一流缓存时早退；测试向两个缓存注入 {0: Mock}（test:271/273）均命中，故不早退，119-121 行随即对 _GLOBAL_COMPUTE_STREAMS 与 _INDEXER_STREAMS 无条件调用 _get_process_cuda_stream。该函数在缓存命中时仍复核 stream.device.index != device_index（attention.py:172）；Mock 的 device.index 为自动生成属性、与整数 0 比较恒不等，于是执行真实 torch.cuda.Stream(device=cuda:0) 构造（attention.py:173）。测试已 patch _cuda_device_index、torch.cuda.current_stream/device/stream/is_current_stream_capturing 等全部入口（test:274-284），但未 patch torch.cuda.Stream 或 _get_process_cuda_stream（该构造在测试文件中仅出现于带 skipUnless 守护的 CUDA 类 364/591 行），所在类 FullOverlapCPUTest（177 行）无任何跳过装饰；无可见 GPU 时第一阶段（capturing=True）即在流规范化处抛异常，用例以 ERROR 结束。
该新增用例在 GPU 隐藏环境确定性 ERROR（测试层面失败，非生产运行时缺陷），污染 CPU-only 套件结果、阻断文档自述的离线验证流程全绿（12 项选定 gate/lifetime 测试中 11 项通过、此项报错）。文档 270-274 行将其表述为 existing/predates the PDL patch，相对本 PR base→head 不成立（测试类与模块均为新增），可能误导维护者按基线遗留问题搁置。生产路径不受影响：真实流对象的 device.index 与缓存键一致，attention.py:172 复核通过即复用缓存，无重建或流churn。
rtp_llm/models_py/modules/dsv41/fp8/decode/full_overlap.py:117： if capturing and any(device not in cache for cache in streams):
rtp_llm/models_py/modules/dsv41/fp8/decode/full_overlap.py:120： _get_process_cuda_stream(cache, x.device) for cache in streams
rtp_llm/models_py/modules/dsv41/fp8/attention.py:172： if stream is None or stream.device.index != device_index:
rtp_llm/models_py/modules/dsv41/fp8/attention.py:173： stream = torch.cuda.Stream(device=torch.device("cuda", device_index))
rtp_llm/models_py/modules/dsv41/fp8/test/test_v41_global_overlap.py:253： def test_cold_capture_and_failed_producer(self):
rtp_llm/models_py/modules/dsv41/fp8/test/test_v41_global_overlap.py:266： main, glob, aux = Mock(), Mock(), Mock()
rtp_llm/models_py/modules/dsv41/fp8/test/test_v41_global_overlap.py:267： glob.cuda_stream, aux.cuda_stream = 123, 456
rtp_llm/models_py/modules/dsv41/fp8/test/test_v41_global_overlap.py:271： full_overlap, "_GLOBAL_COMPUTE_STREAMS", {0: glob}
rtp_llm/models_py/modules/dsv41/fp8/test/test_v41_global_overlap.py:273： full_overlap, "_INDEXER_STREAMS", {0: aux}
rtp_llm/models_py/modules/dsv41/fp8/test/test_v41_global_overlap.py:177：class FullOverlapCPUTest(unittest.TestCase):
rtp_llm/models_py/modules/dsv41/fp8/test/test_v41_global_overlap.py:300：@unittest.skipUnless(torch.cuda.is_available(), "CUDA required")
rtp_llm/models_py/modules/dsv41/fp8/test/test_v41_global_overlap.py:31：class GlobalOverlapCPUTest(unittest.TestCase):
rtp_llm/models_py/modules/dsv41/fp8/decode/global_overlap.py:78： stream = _GLOBAL_COMPUTE_STREAMS.get(device_index)
rtp_llm/models_py/modules/dsv41/fp8/decode/global_overlap.py:87： stream = _GLOBAL_COMPUTE_STREAMS.get(device_index)
docs/benchmark/dsv41_full_layer_optimize.md:270：gate/lifetime tests, 11 passed with GPUs hidden. The existing
docs/benchmark/dsv41_full_layer_optimize.md:273：_get_process_cuda_stream replaces it. It errors without a visible GPU. This
docs/benchmark/dsv41_full_layer_optimize.md:271：FullOverlapCPUTest.test_cold_capture_and_failed_producer still attempts to
rtp_llm/models_py/modules/dsv41/fp8/decode/full_overlap.py:134： global_stream.wait_stream(current)

建议：首选测试侧最小修复：为 glob/aux 设置 device.index=0（Mock 支持链式属性赋值）或 device=torch.device('cuda', 0)，使 attention.py:172 复核命中、Mock 被直接复用，两阶段断言原样保留且全流程 CPU 安全；或在测试中 patch attention._get_process_cuda_stream/torch.cuda.Stream；或将该用例移入带 skipUnless(torch.cuda.is_available()) 的类。生产侧可令 try_full_decode 与 start_global_decode 一致：缓存命中先 .get() 直取、未命中再走 _get_process_cuda_stream，消除姊妹函数分歧。同时把文档 270-274 行 existing/predates the PDL patch 的表述更正为相对上游 base 由本分支引入，并补记 GPU 可见环境下该用例的实测结果。

评审版本：597c930afcb4

rtp-llm-review-bot · 2026-10-20T00:35:59Z

+        # Producer only: QA completion still orders this kernel. The following
+        # QB may initialize early but must GDC-wait before reading Q/scales.
+        tl.inline_asm_elementwise(
+            "griddepcontrol.launch_dependents; // $0",


[P2] FULL decode Q-norm 入口 PDL 触发把 QB 提前启动的读取安全委托给快照外 DeepGEMM 消费端 wait 契约；门仅校验策略位，且构建钉扎 wheel 与文档验证版本不一致

SM10x（capability==10，_v41_decode_global.py:299）、TP1、Index 32×128/TopK512、存在 merged_qb 的 DSv4.1 FULL1/FULL2 解码；部署设置 DSV41_FULL_DECODE_MULTISTREAM=1 与 DSV41_FULL_DECODE_QB_PDL=1，且 DG 全局策略已由部署侧置真（docs:44 列为启用条件，生产代码只读不改，测试断言 set_pdl 未被调用）。eager 预热与 CUDA Graph 捕获/回放均可达（仅冷捕获整体回退，full_overlap.py:117-130）。
base 的该内核无 PDL 分支、query_norm_quant 签名无 pdl 参数（base :56）、compute_qkv.py:98 调用不传 pdl，且 base 全仓 dsv41 无任何 launch_dependents——该边只能等生产者完整完成后隐式释放，文档实测 PDL off 时 norm/QB 零重叠(docs:364)，消费端 wait 位置对本边正确性无关。head 新增入口触发并经 full_overlap.py:172→compute_qkv.py:99 传入，首次打开实测 1.27-4.13us 的提前启动窗口(docs:365)，使外部消费端 wait 成为本边正确性的必要条件。
head 在 _query_norm_quant_kernel 入口、首个 tl.load(:42) 与全部 store(Y:47/Q:57/SF:65) 之前，按 PDL constexpr 发射 griddepcontrol.launch_dependents(:29-39)。当 deep_gemm.get_pdl() 策略为真时，同流紧随的 merged QB（_v41_merged_qb.py:65 直调 deep_gemm.fp8_fp4_gemm_nt，调用处不传任何 pdl 参数，启动属性完全由外部库决定）可以程序化串行属性在 Q-norm 尚未写出 Q/SF 时提前启动；正确性取决于该外部内核（及其 wait 前先启动的辅助内核）在读取 A（量化 Q）/SF（尺度）之前执行 griddepcontrol.wait。生产门 _qb_pdl_enabled(full_overlap.py:23-33) 只读 env 开关与策略位，不校验所装内核的 wait 位置；DeepGEMM 源码不在快照内，仓内无法核实。
若部署 DeepGEMM 在策略为真时确以程序化串行启动 QB，且该内核（或其 wait 前的辅助内核）在读取 A/SF 之前无 griddepcontrol.wait，则 QB 读到未完成的量化 Q 与尺度，Q/index 投影静默数值错误（无异常、无日志），污染该配置下全部 FULL 解码输出；若库遵守自身契约或对该形状未采用程序化启动，则触发仅提前启动时机或为无害空操作。
rtp_llm/models_py/modules/dsv41/_v41_query_quant.py:33："griddepcontrol.launch_dependents; // $0",
rtp_llm/models_py/modules/dsv41/_v41_query_quant.py:42：x = tl.load(X + row * stride_x + cols, cols < D, 0).to(tl.float32)
rtp_llm/models_py/modules/dsv41/_v41_query_quant.py:31：# QB may initialize early but must GDC-wait before reading Q/scales.
rtp_llm/models_py/modules/dsv41/fp8/decode/full_overlap.py:25：os.environ.get("DSV41_FULL_DECODE_QB_PDL", "0") != "1"
rtp_llm/models_py/modules/dsv41/fp8/decode/full_overlap.py:33：return hasattr(deep_gemm, "get_pdl") and deep_gemm.get_pdl()
rtp_llm/models_py/modules/dsv41/fp8/decode/full_overlap.py:172：qb_pdl=_qb_pdl_enabled(mega),
rtp_llm/models_py/modules/dsv41/fp8/decode/compute_qkv.py:99：quantized_qkv = try_project_quantized_qkv(attn, x, pdl=qb_pdl)
rtp_llm/models_py/modules/dsv41/fp8/decode/compute_qkv.py:113：q, index_q = merged.forward_quantized(qr_quantized)
rtp_llm/models_py/modules/dsv41/fp8/_v41_merged_qb.py:65：deep_gemm.fp8_fp4_gemm_nt(
rtp_llm/models_py/modules/dsv41/_v41_query_quant.py:56：def query_norm_quant(x: torch.Tensor, weight: torch.Tensor, eps: float):
rtp_llm/models_py/modules/dsv41/fp8/decode/compute_qkv.py:98：quantized_qkv = try_project_quantized_qkv(attn, x)
docs/benchmark/dsv41_full_layer_optimize.md:364：The norm/QB pair had zero overlap with PDL off and
docs/benchmark/dsv41_full_layer_optimize.md:365：approximately 1.27-4.13 us overlap with the entry trigger, depending on the case.
arch_config/arch_select.bzl:99："deep_gemm@http://artlab.alibaba-inc.com/1/pypi/rtp_llm/deep_gemm/deep_gemm-2.8.2%2Bad62a90.cu132-cp310-cp310-linux_x86_64.whl",
docs/benchmark/dsv41_full_layer_optimize.md:129：2.8.2+30e85d9.cu132, and a FlashMLA wheel exposing Mega Attention.
docs/benchmark/dsv41_full_layer_optimize.md:130：This is an observed environment, not proof that arbitrary mixed wheels work.
3rdparty/deep_gemm/BUILD:17："deep_gemm/include/**/.h",
rtp_llm/test/perf_test/dsv41_full_decode_pdl.py:193：expected = int(cuda.CUgraphDependencyType.CU_GRAPH_DEPENDENCY_TYPE_PROGRAMMATIC)
rtp_llm/models_py/modules/dsv41/fp8/test/test_v41_global_overlap.py:192：dg.set_pdl.assert_not_called()
rtp_llm/models_py/modules/dsv41/fp8/_v41_decode_global.py:299：if torch.cuda.get_device_capability(x.device)[0] != 10:
docs/benchmark/dsv41_full_layer_optimize.md:44：| deep_gemm.get_pdl() | Must already be true for that producer trigger to be enabled |
待确认：部署所用 DeepGEMM wheel（含构建钉扎 2.8.2+ad62a90.cu132）在 get_pdl() 为真时 fp8_fp4_gemm_nt 的启动属性，及其主内核与 wait 前先启动的辅助内核中 griddepcontrol.wait 相对 A/SF 读取的位置——源码不在快照内，该两项前提直接决定所述静默数值错误是否存在

建议：在 _qb_pdl_enabled 或 FULL 预热/首次捕获阶段对所装 DeepGEMM 做完整性审计：仿照仓内 _v41_wo_a_quant.py 的 _VENDOR_HASHES sha256 校验惯例钉扎已验证版本/哈希白名单，或校验 QB 内核 PTX/SASS 中 griddepcontrol.wait 先于 A/SF 装载，不满足则回退为无触发；或首次启用时对 QB 输出与关断 PDL 的参考做逐位自校验（复用 runner compare 逻辑）；同时将构建钉扎 wheel 与文档验证版本对齐并在验证记录中钉扎。

评审版本：597c930afcb4

rtp-llm-review-bot · 2026-10-20T00:36:02Z

+        triton.next_power_of_2(table.shape[1]),
+        num_warps=4,
+    )
+    schedule = _deep_gemm.get_paged_mqa_logits_metadata(


[P2] 一份 DG dense schedule 在一次前向内被多个自源层重复传入 fp8_fp4_paged_mqa_logits，只读契约在快照内无法核实

SM100+DeepGEMM FP4 paged 可用的 DSv4.1 FULL decode：各层 base 每次 fp8_fp4_paged_indexer_score 内新建 request_ids/lengths/tables/schedule get_decode_metadata 仅由一次前向内首个೽ 若 DG 在该调用中改写 schedule/lengths/tables/indices rtp_llm/models_py/modules/dsv41/fp8/_v41_decode_indexer.py:172： rtp_llm/models_py/modules/dsv41/fp8/_v41_decode_indexer.py:180： rtp_llm/models_py/modules/dsv41/fp8/_v41_decode_indexer.py:100： rtp_llm/models_py/modules/dsv41/fp8/_v41_decode_indexer.py:121： rtp_llm/models_py/modules/dsv41/fp8/_indexer_score.py:240： rtp_llm/models_py/modules/dsv41/fp8/_indexer_score.py:247： rtp_llm/models_py/modules/dsv41/fp8/attention_v41.py:2871： rtp_llm/models_py/modules/dsv41/fp8/attention_v41.py:2868： rtp_llm/models_py/modules/dsv41/fp8/decode/rope_metadata.py:27： rtp_llm/models_py/modules/dsv41/fp8/attention_v41.py:2796： rtp_llm/models_py/modules/dsv41/block.py:602： rtp_llm/models_py/modules/dsv41/fp8/test/test_v41_decode_indexer.py:468< rtp_llm/models_py/modules/dsv41/fp8/test/test_v41_decode_indexer.py:570< rtp_llm/models_py/modules/dsv41/fp8/_indexer_score.py:219： rtp_llm/models_py/modules/dsv41/fp8/attention_v41.py:2905：待确认：deep_gemm 包源码不在快照 positions 为 decode_rope_metadata 每前向缓存的同一 converted 张量（默认 DSV41_REUSE_DECODE_ROPE=1），block table 为每步经 bind_attn_cache 绑定给所有层的同一 INDEXER_KV 张量的 [:B] 视图（data_ptr 不变），缓存键命中；默认串行路径每个自源层一次 DG dense 调用，同一 schedule 被逐层多次消费；非默认 DSV41_FULL_DECODE_MULTISTREAM=1 路径同样共享。head 测试明确验证不同 pool 的 owner 共享同一份 metadata。并在同一调用内单次消费（base _indexer_score.py:212-238），score_decode_indexer 无 metadata 参数且调用点不传（base attention_v41.py:2905-2918）；head 新增 get_decode_metadata 缓存与 metadata 旁路，一份 schedule 被一次前向内所有兼容自源层复用，DSV41_FUSE_INDEXER_METADATA 默认开。 8;源层调用 _deep_gemm.get_paged_mqa_logits_metadata 构建 schedule，并与 lengths/safe_table/expanded_table/request_ids 一起按 positions/table 的 data_ptr/shape/stride/dtype 与几何参数缓存（pool 刻意不入键）；FULL 拓扑下其余自源层（layer_id==kv_source_layer_id 且 H32/D128）命中同键，把同一组张量传入各自层的 fp8_fp4_paged_indexer_score→fp8_fp4_paged_mqa_logits。若该 DG 调用改写任一入参，第二个自源层起将消费被改动的调度数据。任一入参：自第二个自源层起调度错乱，logits 行缺失或错位，token 选择静默错误，解码输出质量损坏且无崩溃信号；位于默认开启路径。 schedule = _deep_gemm.get_paged_mqa_logits_metadata( cache[key] = (positions, table, result) os.environ.get("DSV41_FUSE_INDEXER_METADATA", "1") != "1" or not is_supported(positions.device, pool.shape[1]) schedule = metadata.schedule scores = operation( indexer_metadata = decode_indexer.get_decode_metadata( and self.layer_id == self.kv_source_layer_id cache[position_key] = (positions, converted) table = self._block_tables_by_type[INDEXER_KV][:B] with bind_attn_cache(self.attn, kv_cache, block_tables_by_type): /code>： # Owners with identical table views share metadata, never KV pools. /code>： self.assertIs(reused, plan) schedule = _deep_gemm.get_paged_mqa_logits_metadata( all_logits = decode_indexer.score_decode_indexer( 内（deepgemm_wrapper.py 仅从外部包解析符号，仓库内无其实现文件），无法核实 fp8_fp4_paged_mqa_logits 对 schedule_meta/lengths/tables/indices 是否只读
建议：每层打分前单独重建 schedule（其余融合字段仍可共享），或在仓库内补一个多层共享同一 metadata、连续多次真实 DG 消费的严格等值测试（对照 base 稀疏 plan 共享测试的 assertIs+等值模式），以钉住 dense API 只读语义。评审版本：597c930afcb4

rtp-llm-review-bot · 2026-10-20T00:36:06Z

PR #1532 评审：存在待核实的高风险项，等待人工确认

标题：feat: optimize DSv4.1 FULL decode with fused metadata, multistream and minimal PDL（@guluguluhhhh，open）
评审版本：597c930afcb4

待补充的审查依据

DeepGEMM 与 Triton 均为外部依赖，快照无其源码/wheel：QB 消费端 gridDependencySynchronize 相对 A/B 读取的位置、以及 griddepcontrol PTX 在目标工具链的 JIT 编译行为无法直接核验
文档中的历史计时表、32/12 测试计数、SM103 环境版本与六项 CPU-only runner 检查均为运行时观测记录，冻结快照内无法静态复核
deep_gemm(get_paged_mqa_logits_metadata、fp8_fp4_paged_mqa_logits)为外部依赖,快照内无源码,无法核验 schedule 对块长/长度语义及其图重放行为
引擎级 CUDA graph 捕获/回放调度器位于 dsv41 之外未逐行核验(仅依据 _forward_decode_body 每次 forward 调 _begin_forward 的源码事实与测试覆盖)
deep_gemm 库源码不在快照内:FULL 首次将 DG 调用(get_paged_mqa_logits_metadata)置于辅助流并与主流 DG GEMM 并发,库内部是否存在跨流不安全的全局设备状态无法在快照内核实;仅有作者文档检查记录与按位一致测试结果作为旁证
deep_gemm 库源码不在快照内：无法核验 fp8_fp4_gemm_nt 以 PDL 启动时是否在读取激活前执行 griddepcontrol.wait，也无法定位生产侧 set_pdl 调用点
引擎级 CUDA Graph 按 shape 预热/捕获循环位于 dsv41 之外的未改代码，本分片未审阅；FULL 路径各 shape 的 eager 预热是否发生仅能由仓内测试推断
无法执行任何测试或内核(SM10x GPU 不可用)，新增测试是否通过、Triton tl.store 隐式类型转换与 DeepGEMM 调度在图捕获下的实际行为只能静态推断，未能动态验证
快照内没有 torch / cuda-bindings / deep_gemm 源码或版本锁定说明，part-00013 中 CUDAGraph(keep_graph=True)+instantiate()、Event(enable_timing=True, external=True)、cuGraphGetNodes 的返回元数与 CU_GRAPH_KERNEL_NODE_PORT_PROGRAMMATIC、deep_gemm.get_pdl()/set_pdl() 的语义只能按其内部自洽性与文档注释审阅，无法判定在目标运行时是否成立。
按任务约束不执行任何代码，part-00012 中受 SM100 门控的 FullOverlapCudaTest 与 part-00013 全流程只能静态审阅，无法确认捕获/回放的实际数值与时序行为。
torch 2.11+cu130、deep_gemm、flash_mla 与原生 .so 不在快照内：无法确认 torch.cuda.Stream(device=…) 在无 CUDA 环境的行为、Stream.wait_stream 对 Mock 参数的行为、CUDAGraph(keep_graph=…)/Event(external=…) 与 cuda-bindings 图 API 是否存在（影响 CPU 测试失败模式与性能脚本可运行性的最终确认）
deep_gemm 库源不在快照内：get_pdl() 为真时 fp8_fp4_gemm_nt 的等待位置、get_paged_mqa_logits_metadata 在 auxiliary 流并发下的内部状态均无法从快照核验
无执行环境：新增 FullOverlapCPUTest 在无 GPU/有 GPU 两类环境下的失败形态只能由源码与 PyTorch 公开语义推断，未能实跑证实

尚有高风险项待核实

待核实的高风险项

[待核实] QB 入口级 PDL 触发先于全部 Q/SF 写入，读序安全单点依赖快照外 DeepGEMM 消费端读前 gdc.wait，门控仅检查 get_pdl() 且无构建绑定 — rtp_llm/models_py/modules/dsv41/_v41_query_quant.py:33
[待核实] FULL decode Q-norm 入口 PDL 触发把 QB 提前启动的读取安全委托给快照外 DeepGEMM 消费端 wait 契约；门仅校验策略位，且构建钉扎 wheel 与文档验证版本不一致 — rtp_llm/models_py/modules/dsv41/_v41_query_quant.py:33
[待核实] 一份 DG dense schedule 在一次前向内被多个自源层重复传入 fp8_fp4_paged_mqa_logits，只读契约在快照内无法核实 — rtp_llm/models_py/modules/dsv41/fp8/_v41_decode_indexer.py:172

非阻塞发现

[P2] 新增 FullOverlapCPUTest.test_cold_capture_and_failed_producer 在无可见 GPU 环境因 try_full_decode 对 Mock 缓存流重建真实 CUDA 流而确定性 ERROR — rtp_llm/models_py/modules/dsv41/fp8/decode/full_overlap.py:117
[P2] 关闭 DSV41_REUSE_DECODE_ROPE 时,decode_indexer_metadata 缓存逐层插入并钉住整套元数据缓冲,单次前向驻留显存随自源层数线性放大,base 无此保留 — rtp_llm/models_py/modules/dsv41/fp8/decode/rope_metadata.py:9
依据：运维/调试将受支持回退开关 DSV41_REUSE_DECODE_ROPE 置 0 而 DSV41_FUSE_INDEXER_METADATA 保持默认开,SM10x paged 快路径可用(is_supported major==10),FULL 拓扑一次前向含约 59 个 layer_id==kv_source_layer_id 的自源层(docs 记录每次 replay 29 FULL2+30 FULL1;共享拓扑如 [2,8,14,20] 亦为多源),eager decode 或图捕获期间的 Python 执行逐层触发;块表宽度为静态 max_blocks,放大量随部署 max_seq_len 与批量固定成形
base 无 get_decode_metadata/decode_indexer_metadata 缓存(base _select_indices_decode 直接调 score_decode_indexer 且无 metadata 参数);head 新增按 data_ptr 键控、持张量引用、前向级生命周期的缓存,仅当 rope 复用开启(默认)使各层共享同一 converted 且块表同源于共享 metadata 时才退化为单条目
get_decode_metadata 以 positions.data_ptr 为键、以 (positions, table, result) 为值缓存,故意钉住输入防地址复用混叠;DSV41_REUSE_DECODE_ROPE!=1 时 decode_rope_metadata 每层对 int32 的 metadata.position_ids[:T] 执行 .to(torch.long) 新分配 converted(dtype 变必拷贝),上一层 converted 被缓存值钉住、其地址不可被分配器复用,于是每个自源层键必未命中并新插入一整套 lengths/physical_lengths/safe_table/expanded_table(rows×cols int32)/requests/DG schedule;缓存唯一清空点是下一前向最小层 _begin_forward 的整体 pop,前向内无逐出,N 套缓冲并存且跨前向驻留至下一前向开始
单次前向额外驻留约自源层数×(rows×cols×4B+B×cols×4B+小缓冲+schedule) 的显存:文档形状(B64/S6、128K、cols≈1024)单套约 1.8MB、59 层约百 MB 级,更大批量或更长 max_seq_len 下达数百 MB,且保留至下一前向开始才释放,挤占 KV/激活余量,极端配置可触发 OOM;各层元数据内容一致、输出仍正确,属资源耗尽类回归而非正确性类
rtp_llm/models_py/modules/dsv41/fp8/decode/rope_metadata.py:9：if os.environ.get("DSV41_REUSE_DECODE_ROPE", "1") != "1":
rtp_llm/models_py/modules/dsv41/fp8/decode/rope_metadata.py:10：converted = positions.reshape(-1).to(device=table.device, dtype=torch.long)
rtp_llm/models_py/modules/dsv41/fp8/decode/decode_attn_metadata.py:78：position_ids: torch.Tensor # [T_total] int32
rtp_llm/models_py/modules/dsv41/fp8/decode/decode_attn_metadata.py:81：# instead of re-running position_ids.to(torch.long).
rtp_llm/models_py/modules/dsv41/fp8/_v41_decode_indexer.py:128：positions.data_ptr(),
rtp_llm/models_py/modules/dsv41/fp8/_v41_decode_indexer.py:144：cache = shared.setdefault("decode_indexer_metadata", {})
rtp_llm/models_py/modules/dsv41/fp8/_v41_decode_indexer.py:152：expanded_table = torch.empty((rows, table.shape[1]), **options)
rtp_llm/models_py/modules/dsv41/fp8/_v41_decode_indexer.py:180：cache[key] = (positions, table, result)
rtp_llm/models_py/modules/dsv41/fp8/_v41_decode_indexer.py:94：Input references prevent allocator reuse from aliasing a cached key. Pools
rtp_llm/models_py/modules/dsv41/fp8/_v41_decode_indexer.py:100：os.environ.get("DSV41_FUSE_INDEXER_METADATA", "1") != "1"
rtp_llm/models_py/modules/dsv41/fp8/attention_v41.py:1068：if self.layer_id == min(self._shared_attention["layers"]):
rtp_llm/models_py/modules/dsv41/fp8/attention_v41.py:1080：self._shared_attention.pop("decode_indexer_metadata", None)
rtp_llm/models_py/modules/dsv41/fp8/attention_v41.py:3138：self._begin_forward()
rtp_llm/models_py/modules/dsv41/fp8/attention_v41.py:3141：positions, freqs = decode_rope_metadata(
rtp_llm/models_py/modules/dsv41/fp8/attention_v41.py:2868：and self.layer_id == self.kv_source_layer_id
rtp_llm/models_py/modules/dsv41/fp8/attention_v41.py:2871：indexer_metadata = decode_indexer.get_decode_metadata(
rtp_llm/models_py/modules/dsv41/fp8/attention_v41.py:2796：table = self._block_tables_by_type[INDEXER_KV][:B]
rtp_llm/models_py/modules/dsv41/fp8/attention.py:1810：with bind_attn_cache(self, kv_cache, attn_metadata.pool_block_tables):
rtp_llm/models/dsv41_kv_cache.py:82：layer.extend((global_pools[ratio], indexer))
docs/benchmark/dsv41_full_layer_optimize.md:362：Sixteen independent raw traces confirmed 29 FULL2 and 30 FULL1 launches per
rtp_llm/models_py/modules/dsv41/fp8/test/test_v41_decode_indexer.py:623：kv_source_layer_id=layer_id,
rtp_llm/models_py/modules/dsv41/fp8/_indexer_score.py:218：tables = block_table.repeat_interleave(next_n, dim=0).contiguous()
rtp_llm/models_py/modules/dsv41/fp8/_v41_decode_indexer.py:314：safe_table = torch.empty(block_table.shape, dtype=torch.int32, device=q.device)
rtp_llm/models_py/modules/dsv41/fp8/attention_v41.py:2905：all_logits = decode_indexer.score_decode_indexer(
rtp_llm/models_py/modules/dsv41/test/test_v41_decode_rope_metadata.py:34：with patch.dict(os.environ, {"DSV41_REUSE_DECODE_ROPE": "0"}):
建议：缓存有界化:插入前若 positions 指针与既有条目不同则先清空旧条目(维持单条目语义,层内局部引用仍保证在用缓冲不被提前释放);或检测到 DSV41_REUSE_DECODE_ROPE!=1 时 get_decode_metadata 直接返回 None 走每层非融合旧路径;或改用跨层稳定的 position_ids 存储身份作缓存键;并补一条 REUSE=0 与 FUSE=1 组合的驻留回归测试

_{自动代码评审 · head 597c930afcb4 · 存在待核实的高风险项，等待人工确认}

Original-commits: 8e0b3f1

Original-commits: b57181a

Original-commits: 65cedf5, 47b858e, d6826d6, 667f2a1

Original-commits: 468b13b

Original-commits: 1c9d6d0

Original-commits: 58071aa

Original-commits: 914bd96, 105280e

Original-commits: f743725

Original-commits: e5f9c6a, 05fccf2, 801eb19

Original-commits: e13bc06

Original-commits: 8685382, 438fbb2

Original-commits: bdb284c, 3b66ada

Original-commits: 4981fe6

Original-commits: d95fb9c

Original-commits: 8263c6a

Original-commits: d23fb76c

Remove the source-pinned FlashMLA consumer, patch, bindings and outgoing REUSE ready-QKV PDL handshake. Route attention through the existing FlashMLA wheel again. Keep REUSE QKV/index fusion, FULL multistream, metadata fusion and incoming QB PDL intact. Add CPU contracts for the retained producer plan and binding signature.

guluguluhhhh requested review from LLLLKKKK and netaddi as code owners October 19, 2026 14:51

rtp-llm-review-bot added auto-lgtm:needs-approval An unrevoked human CODEOWNER approval is required. review:running Automated review is running; see the progress comment. labels Oct 19, 2026

rtp-llm-review-bot reviewed Oct 20, 2026

View reviewed changes

rtp-llm-review-bot added review:blocked Review found blockers or requires human confirmation. and removed review:running Automated review is running; see the progress comment. labels Oct 20, 2026

guluguluhhhh force-pushed the full_layer_optimize branch from 597c930 to e78fc94 Compare October 20, 2026 05:43

rtp-llm-review-bot added review:running Automated review is running; see the progress comment. and removed review:blocked Review found blockers or requires human confirmation. labels Oct 20, 2026

guluguluhhhh force-pushed the full_layer_optimize branch from e78fc94 to 47b858e Compare October 20, 2026 06:11

rtp-llm-review-bot added review:delayed Review is retrying or has no recent completed task. and removed review:running Automated review is running; see the progress comment. labels Oct 20, 2026

rtp-llm-review-bot removed the review:delayed Review is retrying or has no recent completed task. label Oct 21, 2026

linzehaina force-pushed the full_layer_optimize branch from 8263c6a to b74e9aa Compare October 21, 2026 12:19

rtp-llm-review-bot added review:delayed Review is retrying or has no recent completed task. and removed review:running Automated review is running; see the progress comment. labels Oct 22, 2026

guluguluhhhh and others added 18 commits October 22, 2026 10:49

feat: enable fused V4.1 decode global producer for larger batches

1bc76b6

Original-commits: 8e0b3f1

feat: fuse and share V4.1 FULL decode indexer metadata

03b87e6

Original-commits: b57181a

feat: overlap V4.1 FULL decode attention branches

b6b279e

Original-commits: 65cedf5, 47b858e, d6826d6, 667f2a1

perf: batch speculative masks and reuse complete MTP page tables

7cfcaa0

Original-commits: 468b13b

perf: fuse V4.1 Q-A/KV and BF16 head projection

0cbb044

Original-commits: 1c9d6d0

perf: preserve BF16 sparse logits through token selection

b17aec2

Original-commits: 58071aa

perf: fuse V4.1 QB GEMM with indexer query preparation

6077e0d

Original-commits: 914bd96, 105280e

perf: reuse sparse decode plan lengths across reindex layers

e9a819c

Original-commits: f743725

perf: compact sparse decode plans from candidate flags

ba08a99

Original-commits: e5f9c6a, 05fccf2, 801eb19

perf: add opt-in SM103 merged QB/KV and TopK slot fusion

79aef4a

Original-commits: e13bc06

perf: integrate REUSE QKV index fusion and native FlashMLA PDL

9d31f7f

Original-commits: 8685382, 438fbb2

perf: fuse DSpark MTP preparation and commit row mapping

6df05f7

Original-commits: bdb284c, 3b66ada

refactor: remove deferred sparse plan launch strategy

64ba65c

Original-commits: 4981fe6

fix: flatten CUDA kernel copts selects

85d8a9a

Original-commits: d95fb9c

fix: adapt FULL and Reindex producers to upstream interfaces

efab5b2

Original-commits: 8263c6a

perf: publish compact reindex candidates and graph request validity

7f2a7c9

Original-commits: d23fb76c

build: support locked DeepGEMM device headers

9b18300

guluguluhhhh force-pushed the full_layer_optimize branch from 344d648 to 9b18300 Compare October 22, 2026 02:50

guluguluhhhh changed the title ~~perf: optimize DSv4.1 decode attention, MoE packing, and MTP~~ perf: optimize DSv4.1 decode attention and MTP preparation Oct 22, 2026

rtp-llm-review-bot added review:running Automated review is running; see the progress comment. and removed review:delayed Review is retrying or has no recent completed task. labels Oct 22, 2026

netaddi approved these changes Oct 22, 2026

View reviewed changes

rtp-llm-review-bot added auto-lgtm:authorized Human CODEOWNER authorization exists; current-head review must still pass. and removed auto-lgtm:needs-approval An unrevoked human CODEOWNER approval is required. labels Oct 22, 2026

netaddi merged commit dbfae22 into alibaba:feature/dsv41-dev-main-202161013 Oct 22, 2026
2 checks passed

Control	Default	Behavior
`DSV41_FULL_DECODE_MULTISTREAM`	`1`	Eligible FULL fork/join scheduling
`DSV41_FULL_DECODE_QB_PDL`	`1`	Q-norm producer trigger; also requires eligible merged QB and an already-enabled DeepGEMM PDL policy
`DSV41_FUSE_INDEXER_METADATA`	`1`	Fused, forward-local indexer metadata
`DSV41_FUSED_FRONT_PROJECTION`, `DSV41_FUSED_QB_POSTPROCESS`	`1`	Qualified Reindex projection fusions; `0` restores the corresponding separate operations
`DSV41_OVERLAP_SPARSE_PLAN`	`1`	Asynchronous sparse-plan preparation; `0` uses serial preparation
`DSV41_CUSTOM_QB`, `DSV41_FUSED_QB_KV`, `DSV41_FUSED_TOPK_SLOTS`, `DSV41_FUSED_READY_QKV`	`0`	Optional custom QB/KV/slot/REUSE producers
`DSV4_FUSED_DSPARK_QUERY_INPUTS`, `DSV4_FUSED_DSPARK_COMMIT_ROWS`	`1`	Fused DSpARK proposal/commit preparation

Sunbelt Computer Software

PL/B Language Development and Support

Navigation Menu

Repository navigation

perf: optimize DSv4.1 decode attention and MTP preparation - #1532

perf: optimize DSv4.1 decode attention and MTP preparation#1532
netaddi merged 18 commits into
alibaba:feature/dsv41-dev-main-202161013from
guluguluhhhh:full_layer_optimize

guluguluhhhh commented Oct 19, 2026 •

edited

Loading

Uh oh!

rtp-llm-review-bot commented Oct 19, 2026 •

edited

Loading

Uh oh!

rtp-llm-review-bot Oct 20, 2026

Uh oh!

rtp-llm-review-bot Oct 20, 2026

Uh oh!

rtp-llm-review-bot Oct 20, 2026

Uh oh!

rtp-llm-review-bot Oct 20, 2026

Uh oh!

rtp-llm-review-bot commented Oct 20, 2026

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

7 participants

Sunbelt Computer Software

PL/B Language Development and Support

Conversation

guluguluhhhh commented Oct 19, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Attention Optimizations

Other Decode Optimizations

Configuration

End-to-End Decode Performance

Attention Performance

Recorded Operator Measurements

Accuracy Evaluation

Reproduction and Validation

Uh oh!

rtp-llm-review-bot commented Oct 19, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

自动审查进度

Uh oh!

rtp-llm-review-bot Oct 20, 2026

Choose a reason for hiding this comment

Uh oh!

rtp-llm-review-bot Oct 20, 2026

Choose a reason for hiding this comment

Uh oh!

rtp-llm-review-bot Oct 20, 2026

Choose a reason for hiding this comment

Uh oh!

rtp-llm-review-bot Oct 20, 2026

Choose a reason for hiding this comment

Uh oh!

rtp-llm-review-bot commented Oct 20, 2026

PR #1532 评审：存在待核实的高风险项，等待人工确认

待补充的审查依据

尚有高风险项待核实

待核实的高风险项

非阻塞发现

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

7 participants

guluguluhhhh commented Oct 19, 2026 •

edited

Loading

rtp-llm-review-bot commented Oct 19, 2026 •

edited

Loading