perf: optimize Qwen3.5 MTP decode execution on NPU.

[pjgao]

修改内容与原因

1. 草稿模型主体独立使用 TP1

Qwen3.5 2B MTP草稿模型计算量较小，使用 TP2 时每层通信开销无法被计算有效掩盖。让草稿模型主题不做切分，在各
rank 本地执行避免这部分通信，同时保留体积较大的 LMHead做切分，避免在每张卡上复制完整 LMHead引起大开销。
这样草稿模型主体在每个 target rank 上各自以 TP1 执行，LMHead 仍使用target 的 TP2 group 做分片输出投影。该路径由显式运行时选项控制，并同步调整草稿KV Cache 的容量估算和张量形状。

部分	原来	修改后
Target model body	TP2	TP2
Target LMHead	TP2	TP2
Draft model body	TP2	每个 rank 各自 TP1
Draft embedding	原先可能共享 target embedding	每个 rank 独立加载完整 embedding
Draft LMHead	共享 target TP2 LMHead	独立加载，但仍按 target TP2 group 分片
Draft KV Cache	按 target TP2 shape 分配	按完整 draft heads、world size 1 分配

主模型TP2、草稿模型TP2：

主模型TP2、草稿模型TP1、lm_head TP2：

2. 恢复草稿 forward 异步执行

#1823 (comment) 增加了同步，这里移除每一步草稿 forward 结束后的立即 stream 同步，恢复通过 event 表达依赖的异步执行方式。

逐步同步会把每次草稿计算和后续准备串行化，在 MTP5 下重复形成 host/device 空泡；
event 已能保证真实的数据依赖，不需要在每一步阻塞 host。

3. Qwen3.5 decode projection 直接使用独立投影结果

Qwen3Next 的qkvz/ba 合并权重布局要有split，而Qwen3.5权重本身是分离的qkv、z、a、b projection，无须先拼接再split，#1907 已在prefill/非投机的chunk-prefill阶段修改，这里放开decode/mtp veifify等其余分支限制。

4. Greedy validation 跳过无用 probability

全 greedy MTP 请求直接使用 draft token IDs 和 target token IDs 完成接受判断，不再要求草稿模型返回 probability，也不再构造、拷贝或恢复 draft probability tensor；随机采样路径保持原逻辑。

Greedy 接受判断只比较 token ID，draft probability 不参与计算。跳过这些张量操作可减少每个投机步骤的 softmax、索引和 host/device bookkeeping，同时保持 greedy validation 语义不变。

性能结果

测试条件：Qwen3.5-2B，TP2，MTP5，两张 NPU；开启 ACL Graph，关闭 graph double
buffer 和 prefix cache；max_memory_utilization=0.84。Rank 0 绑定 CPU 240-279，
Rank 1 绑定 CPU 280-319。固定请求集，输出 25 tokens，先 warmup 8 条，再执行
20 条请求 3 轮。以下指标均为 60 条成功请求的平均值。

配置	服务端 E2E (ms)	服务端 TTFT (ms)	服务端 TPOT (ms)	客户端 E2E (ms)	客户端 TTFT (ms)	客户端 TPOT (ms)	接受率
main	186.412	84.150	4.263	191.100	89.437	4.236	62.667%
完整优化	156.252	80.783	3.148	160.426	85.689	3.114	61.579%

完整优化使服务端 E2E 降低 30.160 ms（16.18%），服务端 TPOT 降低
1.115 ms（26.15%）。

各修改收益

每项同时使用按提交顺序累加和从最终代码单独移除的消融实验验证。

下表按修改组顺序累加，每一行均包含该行及其之前的所有修改：

配置	服务端 E2E (ms)	服务端 TTFT (ms)	服务端 TPOT (ms)	客户端 E2E (ms)	客户端 TTFT (ms)	客户端 TPOT (ms)	接受率
main	186.412	84.150	4.263	191.100	89.437	4.236	62.667%
+ 1. 草稿模型主体独立 TP1（LMHead 保持 TP2）	175.780	83.000	3.865	180.541	88.469	3.836	61.579%
+ 2. 恢复草稿 forward 异步执行	162.873	81.767	3.388	167.176	86.784	3.350	61.579%
+ 3. Qwen3.5 decode projection 直接使用独立投影结果	158.552	80.700	3.243	162.991	85.878	3.213	61.579%
+ 4. Greedy validation 跳过无用 probability	156.252	80.783	3.148	160.426	85.689	3.114	61.579%

相邻两行的差值表示对应修改组的累加收益；另外从最终代码中单独移除各组，得到以下
消融贡献：

修改组	累加 E2E 收益 (ms)	单项消融 E2E 贡献 (ms)	单项消融 TPOT 贡献 (ms)
草稿模型主体独立 TP1（LMHead 保持 TP2）	10.632	11.342	0.425
恢复草稿 forward 异步执行	12.907	13.970	0.513
Qwen3.5 decode projection 直接使用独立投影结果	4.321	3.022	0.113
Greedy validation 跳过无用 probability	2.300	3.833	0.132

四组修改在单项验证中均达到 2 ms 服务端 E2E 收益。

验证

• 基于 origin/main@40b5ff57 rebase，递归 submodule 状态一致。
• 使用 MAX_JOBS=16 CTEST_PARALLEL=16 python setup.py build --tilelang-jobs 16 test
  完成 fresh build。
• spec_input_builder_test：22/22 通过。
• sampler_test：13 个通过，2 个仅适用于 MLU 的用例跳过。
• 所有正式性能实验均完成 60/60 请求，prefix cache 始终关闭。