【求助】X99 + 双 7900XTX + SGLang gfx1100 fork：TP2 通了（单流 122–130 t/s），但 RCCL all_gather 有坑、P2P 打不开、aiter 装不上，几点想请教

[南亚阿东]

@Xiaote 感谢，你把我的排查方向从"猜"推到"可排除"。按你的建议做完后，HiCache 那个 SIGABRT 终于抓到了真栈 —— 但结论和"mamba assert"不太一样，另外还排除了三个我自己长期误判的变量。逐条报结果。

---

〇、先认一个错的：之前的结论建立在没法压满的脚本上

我原来的压测脚本用 need_chars = target_tokens * 1.6 估字符数。那是中文系数（约 1.6 字符/token）。换成英文填充之后，同样公式发出去只有 16,720 token，而目标是 55K。

所以之前"ratio 1 的 5/5 全过"那一轮，实际压力只有约 30%，很可能根本没压到 48.6K。这大概是"稳"的来源——不是配置稳，是没压到。

新脚本改成实测回显 ch/tok（我们这套 prompt 实测 5.15–5.34），不再靠猜系数。建议其他人压测时也回显一下，这个坑很安静。

---

一、抓到了真栈：不是 assert，是 GPU 层

启动器加一个环境变量就够了：

export PYTHONFAULTHANDLER=1

崩溃时刻：

rocddevice.cpp:3586: Callback: Queue 0x7b7a1c400000 aborting with error :
    HSA_STATUS_ERROR_EXCEPTION: An HSAIL operation resulted in a hardware
    exception. code: 0x1016
terminate called after throwing an instance of 'c10::AcceleratorError'
  what(): CUDA error: unspecified launch failure   (hipErrorLaunchFailure)
Fatal Python error: Aborted

Can not alloc mamba cache 这个 assert 在崩溃日志里 grep 零命中。

journalctl -k 全历史：ring timeout 0、GPU reset 0、OOM 0 —— 驱动没判定硬件挂死，只是队列报异常。

三组配置，三个不同的崩溃点

配置	崩溃栈顶
HiCache ON + page1	layernorm.py:1164 forward_native ← communicator.py:1343 _gather_hidden_states_and_residual
无 HiCache + page64	hybrid_linear_attn_backend.py:952 _track_mamba_state_extend ← gdn_backend.py:1057
无 HiCache + page1	无 Python 栈，只到 terminate called

三个不同的点，但我判断是同一个底层故障——因为 HIP 文档原文就写了：

CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect.

崩点不是错误点，是 GPU 队列报错时恰好排到的 API。

---

二、四组对照（每轮只改一个变量）

固定：TP2 / gptq / fp8 KV / 196K / mamba bf16 / 40 槽 / 并发 8 / DFLASH。

组	HiCache	AG 补丁	page_size	请求	结果
A	ON ratio1	ON	1	45.9K × 12	第 2 个崩
B	ON ratio1	OFF	1	45.9K × 12	第 2 个崩
C	OFF	ON	1	45.9K × 12	第 6 个崩（累计 229,610）
D	OFF	ON	64	45.9K × 12	第 6 个崩（累计 229,608）

三个结论：

1. 关 HiCache 仍崩 ⇒ HiCache 不是原因，只是让它提前 3 倍到来（第 2 个 vs 第 6 个）。
2. 关掉我自己的 all_gather 补丁仍崩 ⇒ 本地补丁洗清。
3. C 和 D 的累计 token 是 229,610 vs 229,608，崩点都在第 6 个，几乎逐位一致 ⇒ page_size 1 是清白的，那个 +25% 收益不用背这个锅。

---

三、二分长度：不是阈值

长度	每个耗时	数量	累计	结果
7.7K	6.6 s	50	386,750	全过
15.8K	20.5 s	8	110,733	第 8 个崩
23.7K	42 s	4	94,881	全过
31.6K	70 s	3	94,863	全过

两件事：

• 累计量不是原因：386K token 全过，而 229K 就崩。
• 长度也不是确定性阈值：15.8K 崩了，但 23.7K 和 31.6K 更长，没崩。

目前唯一能解释全部现象的是：概率性故障——请求越长，单位时间越可能撞上。15.8K 那次可能只是恰好。但我没做重复测量，所以这个结论只能算推断，不能算定论。 下一步打算同配置跑 3–5 次看崩率。

---

四、关于你 #12 的两条判断

mamba 那条：你说的 assert 确实存在，我在源码里找到三处（mamba_radix_cache.py:1067/1083/1208），其中 1067 是 evict 一次、重试失败就 assert，没有退避。方法（"抓 traceback 看有没有那句 assert"）也是对的——照做之后我才排除掉 mamba。但这次崩溃没走到它。

想请教：PR #36770 那个"跳过这段 radix 缓存、不崩"的改法，具体改在哪一行？ 我想确认它和我看到的 HSA 0x1016 是不是两条完全不同的路径。如果是两条，那我这边的问题可能还有第三层。

另一个可能相关的观察：崩溃前 mamba usage 一直只有 0.10，池远远没满。所以不是池耗尽，更像索引或指针层面的问题。

--hicache-size 那条：我在这个 fork 上复现不出"主机内存 ×2"。hybrid_pool_assembler.py:172 _split_hicache_size 是按各设备池大小比例切分的，不是"同一绝对大小给两份"。实测 ratio 1 得 218,195 tokens / 3.57 GB，跟设备侧 218,194 对得上。

想确认一下：#29037/#32915 描述的是不是"同时设了 --hicache-size 和 --hicache-ratio"的那种组合？ 我们只设了 ratio，可能根本不在那个 bug 的触发面上。

---

五、page_size 那条（顺带回答 Q2）

你给的测法我还没跑完，但已经间接得到一个结论：page1 的长上下文代价不是崩溃（D 组 page64 同样崩）。

崩溃点分散在三处，和 page_size 无关；C/D 组累计量逐位一致这个证据我比较信。

---

六、aiter / marlin

• aiter 确认对 gfx1100 无解：aiter_backend.py 全是 is_gfx95_supported() / ck_tile paged batch_prefill，CDNA 专用。你之前判断一致。
• gptq_marlin 建议按 is_hip() 关掉那条 warning：日志先提示"可以用 gptq_marlin 更快"，真按它改直接 ValueError: gptq_marlin quantization is currently not supported in ROCm。每个新用户都会踩一次。

---

七、想请教的三件事

1. PR #36770 具体改哪一行，以及它和 HSA 0x1016 是不是两条路径。
2. _track_mamba_state_extend 在 GDN 路径上有没有已知的长 prefill 越界。崩溃点在那里，且当时 mamba usage 只有 0.10。track_ssm_h_src 在 GDN 分支是 offset + lens_masked // state_chunk_size，45.9K 请求这个值能到 ~179，我想确认它被索引的那个 buffer 是按什么维度分配的。
3. 谁在 RDNA3 上遇到过 HSA_STATUS_ERROR_EXCEPTION 0x1016。这个码是 MEMORY_APERTURE_VIOLATION，我怀疑是 GPU 侧索引越界写，但没有比 traceback 更深的证据了（AMD_SERIALIZE_KERNEL=3 还没试）。

---

再次感谢 @Xiaote 和 @xiongjie-gao。你那个"先做便宜对照"的建议价值很大——四组对照做完，P2P 和我这边的 HiCache 这两条线都能当场结案，比继续调参省得多。

[Xiaote]

@南亚阿东 四条对照很干净，HiCache/page_size 两条线可以先结案。三个问题分开答。

1）#36770 与 0x1016
PR 改的是新树 unified_cache/components/mamba_component.py：_alloc_mamba_slot→_try_alloc_mamba_slot（约 494 行），assert（约 508 行）换成 if slot is None: self._on_mamba_alloc_failed(); return None；prepare_for_caching_req 三个调用点（约 638/660/674）拿到 None 就 return 0，本 chunk 不缓存、请求保留 live state。

你 f84475c 的对应点是 mamba_radix_cache.py:1061，断言在 1067，调用点在 cache_unfinished_req 740/751/756。不能只删 assert，否则 None 会进 donate_mamba_ping_pong_slot/copy_from；要移植就在 1067 返回 None，并在三个调用点加 if new_slot is None: return _skip_cache_unfinished_req(req)。int8（1083）和 COW（1208）是另两条路径，PR 没覆盖。

这跟 0x1016 是两条路：前者是 host 侧“没有可驱逐槽”的 assert，触发会打 mamba slot alloc failed after eviction；你日志里 grep 不到、usage 0.10，可以划掉。

2）track 索引
h = k.new_empty(B, NT, H, V, K)，varlen 时 B=1，NT = Σ ceil(extend_len_i/64)（FLA CHUNK_SIZE=64），h.squeeze(0) 后是 (NT,H,V,K)。track_ssm_h_src = offset + lens_masked//state_chunk_size，offset 是 (extend_seq_lens-1)//state_chunk_size+1 的 cumsum。只要 state_chunk_size==64，它就是当前 extend 的最后一个 chunk，不越界；45.9K 整段一次 extend 的最大值是 717，不是 179。

179 先对一下来源：如果来自 ~11.4K 的 chunked-prefill 分片，正常；如果 45.9K 整段还是 179，那 state_chunk_size 是 256，和 h 的 64 网格不一致——取错状态，不是越界，但会污染 cache。打印 self.mamba_chunk_size 和 h.shape[0] 即可确认。

你崩在 952 行（final 拷贝）而不是 h gather，更像 dst/slot 生命周期：track_ssm_h_dst/track_ssm_final_dst 来自 ping-pong 的 mamba_track_indices，HiCache 驱逐+并发复用时若 slot 已释放/为 -1/超过 ssm_states.shape[0]，就是 GPU 侧越界写。临时加（别常开）：

assert 0 <= int(track_ssm_h_dst.min()) and int(track_ssm_h_dst.max()) < ssm_states.shape[0]
assert int(track_ssm_h_src.max()) < h.shape[0]

final_src/dst 也查。dst 越界就顺着 mamba_ping_pong_track_buffer/get_mamba_ping_pong_other_idx 查 slot 生命周期，别先怀疑 h。

3）0x1016 往下钉
这个码是 memory aperture violation，但 rocdevice.cpp 的栈只是异步报错点，不是出错 kernel。先 AMD_SERIALIZE_KERNEL=3 HIP_LAUNCH_BLOCKING=1 复跑同一组。公开 issue #41299 是 Qwen3.8-Flash-Next/gfx950 的 0x1016，serial 后能过，更像 race/缺同步；#43730（gfx942，~50K prompt）也是长度相关、崩点漂移。你这组“三个崩点、page1/page64 都崩、累计 token 逐位一致”很像同一类。

再二分：--disable-overlap-schedule、--disable-cuda-graph（或只关 piecewise）、--enable-mixed-chunk 关掉，各跑 45.9K×12。mixed-chunk + extra_buffer 有已知 bug（#39342：混合批跳过 checkpoint 写但仍 donate slot，fix PR #39526 还开着），不直接造成 0x1016 也建议先关掉排除。

公开面没找到 gfx1100 的 0x1016 案例；RDNA3 支持还在 tracking #30599，你这台可能真是首批。serialize 后还崩，把第一份失败栈和 dmesg | grep -i amdgpu 发我，我按 kernel 名对。

另外两条翻案我认：--hicache-size 的 2× 是 #32915 之前的老树，你的 _split_hicache_size 已按设备池比例切总预算；只设 ratio 根本不走 size split，这条划掉。mamba assert 也是我上轮按旧代码给的优先级，实际没走到，抱歉。

[南亚阿东]

@Xiaote 按你给的三步做了，先报结果：serialize 后仍崩，断言方向(dst/slot)和 h_src 都没越界。另外我上一帖里有个数字报错了，先更正。

---

〇、更正：我算的 179 是错的，你算的 717 才对

你指出的两点我都核实了：

self._mamba_chunk_size = getattr(hf_text_config, "mamba_chunk_size", 64)

我们的 config.json 里没有 mamba_chunk_size 字段，所以取默认 64。实测打印也是 chunk=64。和 FLA 的网格一致，不是 256。 179 是我拿 mamba_track_interval(256) 去算的，用错了常量。

不过你说的"整段 45.9K 最大值 717"在我们这里是 128，原因是 --chunked-prefill-size 8192：每个 forward 只吃 8192 token，NT = 8192/64 = 128。实测：

[AGU] h.shape=(1, 128, 24, 128, 128) dim=5 contig=True
       → (B, NT, H, V, K)
[AGU] ext_lens=[8192] prefix=[0] track_seqlens=[8192]

---

一、serialize 复跑：仍崩，但崩点稳定了

export AMD_SERIALIZE_KERNEL=3
export HIP_LAUNCH_BLOCKING=1

组	配置	崩溃位置	累计 token
基线	无 serialize + page1	req#6	229,610
G1	serialize + page1	req#4	137,766

还在崩，所以不是纯 race。 但 serialize 让它提前两个请求崩，且栈固定在 _track_mamba_state_extend:952（以前是漂的）。

我理解这是：同步执行后，出错点不再被异步推迟到后面的 API，所以报告点 = 出错点。按这个读，更像是确定性的越界写，而不是时序问题。

---

二、断言结果：两个方向都没越界

按你给的断言（外加我多查的 final_src / recompute_dst），实际数据：

[AGU] ssm_states.shape = (41, 24, 128, 128)     # 41 个槽
[AGU] h_src  n=2 vals=[77, 79] max=79           # NT=128  → 77,79 合法
[AGU] h_dst  n=2 vals=[11, 16] max=16           # 槽 41    → 11,16 合法
[AGU] final_src n=1 vals=[2] max=2
[AGU] final_dst n=1 vals=[3] max=3

全部在范围内。 你怀疑的 dst/slot 生命周期，我顺着看了 h_dst=[11,16] 和 final_dst=[3]，没有 -1、没有超 41。

所以这条也要划掉——但我想确认一件事：你给的 h.shape[0] 检查，在我的 5 维布局下应该是第 2 维（NT）。我一开始按 h.shape[0]=1 判，差点得出"77 越界"的假结论。是不是你那边 h 是 (NT,H,V,K)（被 squeeze 过），所以用 shape[0] 是对的？

---

三、这轮唯一的新观察：崩溃前是2 条序列合批

崩溃前的 prefill 日志（serialize 模式）：

#new-seq: 1, #new-token: 8192, full token usage: 0.21, mamba usage: 0.17  ×5
#new-seq: 2, #new-token: 5084, full token usage: 0.21, mamba usage: 0.20   ← 崩在这

最后一个不完整 chunk（5084 token）变成了 #new-seq: 2，而且 h_src/h_dst 正好都是 n=2。

这和 _init_track_ssm_indices 的 GDN 分支对得上：

num_h_states = (extend_seq_lens - 1) // state_chunk_size + 1
track_ssm_src_offset[1:] = torch.cumsum(num_h_states[:-1], dim=0)
track_ssm_h_src = track_ssm_src_offset[mamba_track_mask][not_aligned] \
                + (lens_masked[not_aligned] // state_chunk_size)

src_offset 是按批内前序序列累加的，所以第 2 条序列的 offset 会是第一条的 chunk 数。如果两条序列共享同一个 h 但 offset 没考虑合批，h_src 就会指到错误的位置——这次数值恰好没越界（79 < 128），但很可能取错了 chunk。

想请教：#new-seq: 2 是什么在合批？ 是 chunked-prefill 的两段合并调度，还是 DFLASH 主/草稿模型？

• 如果是 chunked-prefill 合并，那 src_offset 这段的 extend_seq_lens 到底对应哪一层，我想对着看一遍；
• 如果是 DFLASH 双模型，那 h 的 B 维就可能是 2 而不是 1，我现在的 (1,128,...) 只抓到了一次，需要再抓。

---

四、你说的另外几条

① #36770 的移植点我核对过了，你说的位置对得上：

• 我们 fork 的 mamba_radix_cache.py:1061 是 _alloc_mamba_slot，断言在 1067
• 调用点在 cache_unfinished_req 的 740/751/756
• 你提醒"不能只删 assert，否则 None 会进 donate_mamba_ping_pong_slot/copy_from"——这条我记下了，但这次崩溃不走这条路，所以先不移植。int8(1083)/COW(1208) 那两条也没覆盖，一并记着。

② --disable-overlap-schedule：还没跑，下一步做（既然 serialize 都不救，想先看它是否改变崩点位置）。

③ mixed-chunk：我们服务里 enable_mixed_chunk 本来就是 False，这条不适用。

④ --disable-cuda-graph：还没试。

⑤ 公开 issue 我读了：

• #41299（gfx950，serial 后能过）——我们 serial 后还是崩，所以和它不是同一类；
• #43730（gfx942，~50K prompt，崩点漂移）——症状最像，我正在比对它的复现条件。

⑥ gfx1100 无公开案例：理解，那就当首批处理。

---

五、我这边的一个教训（可能对别人有用）

这一轮我连续报了三次错误的"越界证据"，全是因为读错维度：

1. 用 mamba_track_interval=256 算 chunk 序号（应该用 64）
2. 把 h.shape=(1,128,24,128,128) 前两维读成 (2,24)
3. 用 h.shape[1] 当行数上限，但真实行维是 h.shape[1] 在 squeeze 后才成立

结论：验证 track_ssm_h_src 上限前，必须先打印 h.dim() 和完整 h.shape，不能凭 h.shape[x] 猜。 我把这段写进了自己的排查记录。

[Xiaote]

@南亚阿东 两个问题分开答，先把维度这条认了。

1）h.shape 那条是我给错了
5 维 h=(B,NT,H,V,K) 里 varlen 的 B 恒为 1，真正被 track_ssm_h_src 索引的是 NT。_track_mamba_state_extend 里先 h = h.squeeze(0) 再 h[src]，所以函数内用 h.shape[0] 是对的；你在外面打印 5 维张量时应该用 h.shape[1]。我上轮写 h.shape[0] 是按 squeeze 后的视图说的，没写清楚，抱歉。

2）#new-seq:2 是两个请求合批，不是 chunk 合并，也不是 DFLASH
#new-seq 就是 PrefillAdder.can_run_list 的长度，即这一轮从 waiting queue 新放进来几个请求。你并发 8，--chunked-prefill-size 8192 下前一个请求的尾 chunk 只占 5084，剩下的预算会再塞一个等待请求进来，于是 #new-seq:2。extend_seq_lens 是这一轮每个请求各处理多少新 token 的向量，长度就是 batch 里的请求数；h 是 (1, Σceil(extend_len_i/64), H,V,K)，按请求顺序摊平，不是 (2,...)。
DFLASH 是投机解码，draft/verify 都在同一个请求里跑，不会给 scheduler 增加 seq；prefill 的 draft 前向也不会进 #new-seq。

3）你的数其实正好对上两条序列，79 不越界
#new-token:5084 + h_src=[77,79] 可以拆成两条都不整除 64 的序列，例如 4957 + 127：

• seq0：ceil(4957/64)=78，4957//64=77 → h_src=77
• seq1：offset=78，127//64=1 → h_src=79
• NT=78+2=80，79 是最后一行，合法。

你那个 h.shape=(1,128,...) 是 8192 整 chunk 的日志，不是崩的那一轮；崩的那轮 NT 大概率是 80。把崩溃 forward 的 h.shape[1]、extend_seq_lens、extend_prefix_lens、mamba_track_seqlens、mamba_track_mask 一起打出来。如果 h.shape[1]==Σceil(extend/64)，这段 offset 就是自洽的，先别再往越界查。

4）下一步建议按这个顺序
① 在 _track_mamba_state_extend 的 copy 前后各加一次 torch.cuda.synchronize() 并打点：若 "before copy" 没打出来就崩，错在 extend kernel；打出来后才崩，才是 copy 本身。
② 客户端改成并发 1 跑同一组 45.9K×12。serialize 会让 prefill 变慢，若你是固定速率发请求，队列深度和合批组成会变，这不是单变量对比。并发 1 还崩 → 单序列尾 chunk；不崩 → 合批路径。
③ 再试 --disable-overlap-schedule，以及把 prompt 补齐到 64 的倍数（如 49152）对比 45.9K 的非对齐尾，看是不是尾 chunk 触发。
④ 抓最后 kernel：rocprofv3 --kernel-trace -o trace -- python -m sglang.launch_server ...，0x1016 的 Python 栈经常是报错点不是出错 kernel。

如果 ① 证明是 copy 本身，再查 ssm_states 是否 strided/alias（打印 is_contiguous()、stride()、storage_offset()）；索引都在范围内的话，剩下更可能是源行没被写或目标视图 stride 不对，而不是 h_src 算错。

[南亚阿东]

@Xiaote ① ② 都做完了，②（并发 1）直接把你说的"合批"坐实了。另外想请教一个和主板 BIOS / PLX 窗口相关的底层问题，方向可能和我之前想的不一样。

---

〇、先谢两条纠正

① h.shape 那条：懂了，是我们打印位置不同。函数内 h = h.squeeze(0) 之后用 shape[0] 是对的，我在外面打 5 维张量该用 shape[1]。两边都没错，是我没意识到 squeeze。

② #new-seq: 2 是合批：你那句"前一个请求尾 chunk 只占 5084，剩下的预算会再塞一个等待请求"，我拿并发 1 验证了 —— 完全正确。

③ 我拿 h.shape[1] 之前抓到的是 8192 整 chunk 的帧，你说得对，不是崩的那一轮。后面我会按你的要求把崩溃 forward 的五个量一起打出来。

你那个 4957 + 127 的拆解我验算了一遍，和我们的数据完全对得上：

seq0: ceil(4957/64)=78，4957//64=77   → h_src = 0 + 77 = 77
seq1: offset = 78，127//64=1          → h_src = 78 + 1 = 79
NT = 78 + 2 = 80，79 = 最后一行，合法

⇒ h_src=[77,79] 在 NT=80 下完全自洽，越界这条可以彻底划掉。

---

一、决定性对照：并发 1 vs 并发 8

按你的建议改了客户端并发（不是改服务端），其他全不变（page1 / 无 HiCache / DFLASH / 40 槽）：

组	并发	每请求实测	请求数	累计 token	结果
基线	8	45.9K	12	229,608	崩于第 6 个
并发 1	1	53.1K	12	636,710	12/12 全过

明细：

req#1..12 全部 OK，每个 ~184s，prompt=53059/53060
累计 636,710 token，服务存活

并发 1 的请求更长（53K vs 46K）、累计量接近 3 倍，却一个都没崩。

⇒ 合批路径是触发条件，单序列长 prefill 完全没问题。

这也解释了我之前那组"7.7K × 50 全过"：7.7K 被切成 8192 整 chunk，尾 chunk 几乎不留预算，很少能合成 #new-seq: 2。而 45.9K 的最后一个 chunk 正好剩 5084，空间刚好够塞第二个。

另外谢谢你指出 serialize 那组的对比问题——序列化会改变 prefill 速度，固定速率发请求时队列深度和合批组成都会变。所以 G1 那个"第 4 个 vs 第 6 个"我改成定性结论处理，不再当定量数据。

---

二、下一步我打算查合批路径里的哪几处

既然锁定了合批，我准备按你 #16 的 ① 做，但专门打 #new-seq: 2 那一轮：

# _track_mamba_state_extend 的 copy 前后
print("[BEFORE] h.shape=%s ext_lens=%s prefix=%s track_seqlens=%s mask=%s"
      % (h.shape, forward_batch.extend_seq_lens.tolist(),
         forward_batch.extend_prefix_lens.tolist(),
         forward_batch.mamba_track_seqlens.tolist(),
         forward_batch.mamba_track_mask.tolist()), flush=True)
torch.cuda.synchronize()
print("[AFTER-SYNC]", flush=True)

判据按你说的："BEFORE" 没打出来 ⇒ 错在 extend kernel；打出来才崩 ⇒ 才是我后面这几行 copy。

---

三、想请教一个方向（和 0x1016 无关，是 PLX / BIOS 那一层）

这条本来是我自己想了很久但一直没敢动的，读到你的回复后觉得也许该正经问一下。

背景：我现在是"跨 PLX 双 32G BAR"

这台 X99-E WS（BIOS 4001 原厂，非 USB3.1 版）现在两卡在 1 槽 + 4 槽，实测：

07:00.0 (1 槽, PLX#1)  Region 0: 32G @ 381000000000
0d:00.0 (4 槽, PLX#2)  Region 0: 32G @ 380000000000

04:10.0 (PLX#1 上游)  Prefetchable: 380c00000000-3817ffffffff  [size=48G]
0a:10.0 (PLX#2 上游)  Prefetchable: 380000000000-380bffffffff  [size=48G]
04:08.0 / 0a:08.0     Prefetchable: [disabled]

每颗 PLX 的 upstream prefetchable window 是 48G。

所以如果两卡塞进同一个 PLX，需要 2 × 32G = 64G > 48G，BAR 必然掉 256M —— 这和我 10-31 实测"1+3 槽 BAR 掉 256M"一致（1、2、3 槽同属 PLX#1）。

但我在 Win-Raid 看到一条矛盾，想请你判一下

同一个论坛帖（Asus X99-E WS/USB3.1 BIOS Mods incl. ReBar Support）里：

• 有人提供改装 BIOS：插入 ReBarDxe v0.3、MMCFG Base 1GiB→2GiB、MMIO High Size 256GiB→1TiB，并有人用它在 Above 4G 下插 6 块 Tesla P40（每块 32G BAR）
• 另一位用户（ubr）报告：用 PCIEX16_1 + PCIEX16_3 两个槽，2× 32G BAR 卡跑通了
• 但同时另一位（becca）刷了同一个改装 BIOS，/proc/iomem 仍是 256GB，PLX 口 window 没变，第二个卡没地方枚举

⇒ 同一块板、同一个 BIOS 文件，有人成功有人失败。1+3 槽到底是"同 PLX 必掉"还是"BIOS 拉开窗口后可行"，我这边没有定论。

三个具体问题

① PLX PEX8747 的 upstream prefetchable window，是 POST 时 BIOS 静态分配的，还是可以被后续 resize 的？
我的理解是 PEX8747 是 non-transparent 之外的普通 PCIe switch，它的 bridge window 由 BIOS 在 bus enumeration 阶段写死；Linux 的 pci=realloc 只能重排，不能超出 BIOS 给的父窗口。这个理解对吗？

② 那「MMIO High Size 256GiB→1TiB」改的到底是哪一层？
是给 CPU root complex 的 MMIO 上限、还是给 PLX 这类 bridge 的 window 预算？如果只改前者，PLX 自己的 48G 是 PLX 固件/BIOS 里单独设定的，那这个 patch 对"同 PLX 双 32G"就没用——和 becca 的结果一致。

③ 如果要让"同 PLX 双 32G"成立，责任在哪一层？

• BIOS 层：能不能通过 RU.EFI 之类直接改 PLX 上游 window 的大小变量？
• 内核层：你之前提的 PCI: Reserve prefetchable window headroom for Resizable BARs（series 520756） 能否解决"bridge 初始 window 太小"？还是它只管"设备 resize 时预留余量"，仍然受限于 bridge 自己的 max window？
• 还是说 在 E WS 这种 PLX 双跳拓扑上，这条路根本不通，只能维持跨 PLX？

补充一个我实测到的、可能相关的数字

同一份配置，只改 BAR 大小（32G vs 256M）：

	32 GB BAR	256 MB BAR	损失
单流	149.5	132.5	−11.6%
4 路	307.8	276.2	−10.3%

⇒ ReBAR 在我们这个负载上值 11.6%，所以如果"同 PLX + 双 32G + P2P"能成，它的收益就不是我上次估的 +1~2%，而是 11.6% + P2P 的额外收益。这改变了我对升内核 / 改 BIOS 的性价比判断。

---

四、如果你觉得这条路不值得走

也请直说。我现在的判断是：

• BIOS 改动风险：第三方 Mod BIOS，刷废要编程器救砖；而且 becca 的先例说明刷了也不一定生效
• 内核补丁风险：要上 7.0，ROCm 7.2.3 不在官方矩阵里
• 确定收益：11.6%（BAR）+ 未知（P2P）

如果结论是"维持跨 PLX、别动 BIOS"，我就不折腾了，把精力放回合批这个问题上。

再次感谢，你这两轮把 P2P 和崩溃两条线的排查范围都收得很干净。

[Xiaote]

@南亚阿东 三个问题按层次答，结论先说：瓶颈在 BIOS 枚举，不在 PEX8747 硬件；但为这 11.6% 刷 Mod BIOS 不值——跨 PLX 双 32G 已经把这 11.6% 拿到了。

① PEX8747 的 window 是 BIOS 写的，不是硬件上限
它是普通 PCI-to-PCI bridge，上游口 Prefetchable Base/Limit 是可写寄存器，POST 时由 BIOS 枚举写入；Linux 用 pci=realloc 也能重写、必要时能扩大，不只是“重排”。上限只受父窗口（root complex 的 MMIO aperture）约束。48G 是这次 BIOS 给的值，不是 PEX8747 的硬件上限。

② MMIO High Size 改的是父预算
它设的是 CPU root complex 在 4G 以上的 MMIO aperture，BIOS 再拿这个预算分给各 bridge。256G→1T 只是父窗口变大；BIOS 若仍按“每 PLX 48G”分，PLX 上游不会自动变 64G。becca 刷完 /proc/iomem 仍是 256G，说明那版 mod 的 1T 没生效（或被这版 BIOS 忽略），不能反证父窗口没用。

③ 责任层与补丁

• BIOS 是决定者：上下游窗口都在枚举时写死；持久改只能改 BIOS 的资源分配。RU.EFI 只能当场改寄存器，重启失效。
• 内核：pci=realloc 能在父窗口内重建（理论可把 04:10.0 拉到 64G），但要 64G 对齐空洞，还可能移动启动盘资源——你 6.8 上卡 30 分钟就是这个。
• headroom 补丁（LKML v3）：补的是“BIOS 按 boot-time 小 BAR 分小了，驱动后 resize 没空间”，做法是 sizing 时按 ReBAR max 预留、再释放重算 pref window。它变不出父窗口，且只在 max > current 时触发。你 7900 XTX 启动 256M、max 32G 正好命中，但两卡要 64G，仍看 root 有没有 64G 空洞，还要上 7.0+补丁+pci=realloc，不是打上就成。

Win-Raid 那个矛盾：ubr 1+3 插 2×32G V100 成功，说明 mod BIOS 不是做不到；但 V100 的 32G BAR1 是原生上报，你的 7900 XTX 默认 256M、靠 OS resize——BIOS 枚举时看不到 32G 需求。这正是补丁想补的缺口，两者不能直接照搬。

建议：这条线先封存

1. 同 PLX 不会再多 11.6%——你比的是 32G vs 256M，跨 PLX 已是 32G；同 PLX 只多“可能 P2P”，还可能因 64G 塞不进 48G 反掉一张到 256M。
2. 四组拓扑 can_access_peer 全 False，P2P 卡在更上层（ACS/IOMMU/root complex），不是同 PLX 就自动通。先查 lspci -vvv -s 04:10.0 与下游口 ACS、dmesg | grep -i iommu，试一次 amd_iommu=off iommu=pt 看 rocm-smi --showtopo。比刷 BIOS 便宜。
3. 真想验同 PLX，先别刷 BIOS：两卡插 1+3，看 04:10.0 的 window 是否 ≥64G、第二张卡是否仍 32G；不行插回 1+4，十分钟可逆。
4. 现在真正卡你的是合批 0x1016 和 RCCL all_gather 原生 bug，收益比 11.6% 大。合批按计划打 BEFORE/AFTER-SYNC；重点看 #new-seq:2 时两条序列的 extend_prefix_lens / mamba_track_mask。

[南亚阿东]

@Xiaote @xiongjie-gao 又是同一题（topic 2774 原帖），但这次我想换个问法：不再报猜测的机制，只报一组数据和一个排除结果。

先把两个之前的错误更正，再报崩率。

---

〇、更正：我之前说"serialize 后仍崩"，那个实验是无效的

我按你 #14 的建议加了 AMD_SERIALIZE_KERNEL=3 + HIP_LAUNCH_BLOCKING=1 复跑，报"仍崩 ⇒ 不是纯 race"。

后来发现那次变量根本没生效。

判据很简单：那组每请求耗时 140.3 / 141.4 / 141.8 s，和不 serialize 的 137–140 s 完全一样。如果真序列化，prefill 的几千次 kernel 发射全部同步，不可能只差 1 秒。

然后我又犯了一个错：第二次做 blocking 对照时，启动脚本写了 if ! alive 就跳过启动，结果复用了上一个实验遗留的服务，HIP_LAUNCH_BLOCKING 连进程环境里都没有。

⇒ 你给的第一条建议，我执行了两轮都是错的。抱歉。

并且我确认了这个变量在本机是有效的（这一点很重要，见下）：

500 次 256x256 matmul
default              33.0 ms
HIP_LAUNCH_BLOCKING  90.0 ms   ← 2.7x

教训：判断环境变量是否生效，不能只看 /proc/<pid>/environ（有 ≠ 生效），要看可观测的性能差异。

---

一、崩率：并发 8 + 45.9K，5 次里 2 次崩

配置固定不变（page1 / 无 HiCache / DFLASH / 40 槽 / 并发 8），同一台机器、同一份代码，每次 12 个 45.9K 请求：

run	结果	崩在第几个	累计 token
1	崩	第 3 个	91,843
2	活	—	551,067
3	活	—	551,065
4	活	—	551,065
5	崩	第 1 个	0

崩率 2/5 = 40%，崩点从第 1 个漂移到第 3 个。

三次成功时 token 逐位一致（551,067 / 551,065 / 551,065，耗时 1674–1678 s）——
一旦过了崩溃窗口，系统就完全正常。

把之前所有记录摊开：

记录	并发	请求	结果
基线 ×2	8	45.9K	第 6 个
崩率 run1	8	45.9K	第 3 个
崩率 run5	8	45.9K	第 1 个
K2(blocking)	8	45.9K	第 4 个
崩率 run2/3/4	8	45.9K	12/12 ×3
探针组	8	45.9K	11/11
并发 1	1	53K	12/12 / 636K
7.7K 组	8	7.7K	50/50 / 386K

并发 8 + 45.9K：8 次里 4 次崩（≈50%）。并发 1：3 次全过，累计 636K。

这组数字排除了什么

任何确定性的配置错误都不会呈现这个形态。 同配置有时第 1 个就崩、有时 12 个全过——
所以"参数调错了""buffer 算小了""版本不兼容"这一类都可以排除。

---

二、用你建议的 serialize 判定，结论是：不是 GPU kernel 竞态

修好启动后重做（HIP_LAUNCH_BLOCKING=1，已用 matmul 验证生效 2.7x）：

组	每请求耗时	结果
普通	140 s	崩率 40%
blocking 生效	141 s	崩于第 4 个

prefill 只慢 1% —— 说明这个负载对 kernel 发射同步不敏感（大部分时间在一次发射内的 GPU 计算里），
但变量本身确实生效。

⇒ blocking 真正开启的情况下仍然崩，可以排除「GPU kernel 之间的异步执行竞态」。

---

三、合批那条也不是（我之前推得太急）

你 #16 说 #new-seq: 2 是尾 chunk + 新请求合批，我拿并发 1 验证了它的存在性是对的。
但我后来把它当成了触发条件，下面这个反例推翻了：

崩率 run1 崩于第 3 个，但该次日志里 #new-seq: 2 出现 0 次（全是 #new-seq: 1）。

而探针组（专打 nseq>=2）一次合批都没有，却 11 个全过。

⇒ 合批既非充分也非必要。

---

四、现在手上的排除清单（每条都有对照）

假设	结论	判据
mamba assert（#36770 那条）	不是	崩溃日志 grep 零命中
HiCache	不是	关 HiCache 仍崩
本地 AG 补丁	不是	关掉仍崩
page_size 1	不是	page64 仍崩，累计 token 逐位一致（229,610 vs 229,608）
请求长度 / 累计量	不是	并发 1 时 636K 全过
h_src / h_dst 越界	不是	h_src=[77,79] < NT=80、h_dst=[11,16] < 41，全部合法
合批路径	不是	见 §三
GPU kernel 异步竞态	不是	HIP_LAUNCH_BLOCKING 生效后仍崩
ring timeout / GPU reset / OOM	全 0	journalctl -k

---

五、想请教的三件事

① 有没有 concurrent-only、崩点漂移的 HSA 0x1016 先例？

你提过的 #43730（gfx942，~50K prompt，崩点漂移） 症状最像，我读了但没法确认它是否也是并发相关。
如果它是单请求复现，那和我们不是一类；如果也是"并发才崩"，那可能值得顺着它的 issue 线索查。

② blocking 排除了 GPU kernel 竞态之后，下一个该排除什么？

我现在的候选，按成本排：

• --disable-overlap-schedule（这是唯一还没测的"CPU 侧异步"开关）
• --chunked-prefill-size 16384（减少 chunk 次数，看崩率变不变）
• 并发 2 / 4 / 6 各 3 次（看崩率是否随并发单调）
• rocprofv3 --kernel-trace（你 #16 ④ 提的，还没做）

你觉得哪个最该先做？ 我个人倾向 --disable-overlap-schedule，因为 blocking 只作用于 GPU 发射，
而 SGLang 的 overlap schedule 是 CPU 侧的——它才是"并发才出问题"最可能的原因。

③ 有没有办法让 0x1016 报出真正的出错 kernel？

Python 栈三处不同（layernorm.forward_native / _track_mamba_state_extend / 无栈），
按你之前的说法这是异步报错点。AMD_SERIALIZE_KERNEL=3 我怀疑也没生效（和 blocking 同一个坑），
rocprofv3 是唯一能拿到 kernel 级的办法吗？

---

目前我不敢再给机制解释了——这一轮我已经推错五六个方向，每次都是被后续数据推翻。
所以这次只报数字，问题也尽量问在"该排除什么"和"哪里能拿到更硬的证据"上。

再次感谢，你那两条（"先做便宜对照"、"serialize 先跑"）到现在仍然是最有价值的方向，只是我执行时踩了两个坑。

[Xiaote]

@南亚阿东 这轮数据够硬，先按“排除”答，不再给机制。

1）#43730 不是并发类，别顺着它查并发
我翻了原 issue（sgl-project/sglang#43730）：8×MI300X gfx942、GLM-5.3、NGRAM 投机，单请求 ~50K prompt；prefill 完成、HTTP 200 开始后，第一个 decode step 所有 rank 一起 0x1016。同一台机器关掉 spec 能跑到 1M；chunked-prefill 32K/131K、page-size 32 都崩，是长度相关。issue 下 0 评论，作者也还没试 serialize。
所以它和你“并发 1 + DFLASH 53K 12/12 过”不是同一触发条件。但它是你还没做过的干净对照：关 spec。 去掉 --speculative-algorithm DFLASH 和 draft 参数，其余不变，并发 8 + 45.9K 跑 5 次。仍崩 → spec 划掉；不崩 → 再拆 DFLASH 和显存布局（比如 --speculative-num-draft-tokens 1 或换算法），别直接判 DFLASH 有 bug。

2）下一个开关：先关 spec，再 --disable-overlap-schedule
你的直觉方向对，但 blocking 那条排除要收窄：HIP_LAUNCH_BLOCKING=1 / AMD_SERIALIZE_KERNEL 只序列化普通 kernel enqueue，不覆盖 hipMemcpyAsync 和 CUDA graph replay。你开了 --cuda-graph-bs-decode 1..8，并发 8 时 decode graph 完全可能和另一个请求的 prefill 重叠；_track_mamba_state_extend 崩在 prefill，这个组合正好没被 blocking 排掉。
建议顺序：
① 关 spec，5 次；
② 加 --disable-overlap-schedule，5 次；
③ 还崩再关 decode graph（--cuda-graph-bs-decode 置空；fork 若支持 --disable-cuda-graph 也行）跑一轮。
--chunked-prefill-size 16384 和并发 2/4/6 扫描先放后面：前者只改合批概率，不能排除假设；后者 3×3 次太贵，等前两个开关出结果再扫。

3）拿真正出错的 kernel
rocprofv3 --kernel-trace 只能看时间线，异步 0x1016 下“最后完成的 kernel”不一定是出错的那个。更硬的是 rocgdb 的 precise memory，它会在 GPU 访问违例的 wave/指令处停：

rocgdb --args python -m sglang.launch_server ...
(gdb) set pagination off
(gdb) set amdgpu precise-memory
(gdb) catch signal SIGSEGV
(gdb) catch signal SIGABRT
(gdb) catch signal SIGBUS
(gdb) run
# 崩后
(gdb) thread apply all bt full
(gdb) info queues
(gdb) info dispatches

只需复现一次，慢没关系。另外崩完立刻 dmesg -T | grep -iE 'amdgpu|kfd|fault|VM_L2|IH'；ring timeout 0 不等于没有 GPU page fault，faulting VA/pasid 能直接指到 buffer。
还有一点：你只开了 kernel serialize，copy 没管；补一轮 AMD_SERIALIZE_COPY=3 + AMD_SERIALIZE_KERNEL=3 一起开，才能把“copy 与 kernel 之间”的窗口也关掉。变量生效别再看耗时 1%，用 rocprofv3 看是否还有重叠。

先做 ① + copy serialize，再上 ②；这两步能把嫌疑收成“spec 路径 / graph+overlap / 真越界”三类，再上 rocgdb 才不浪费。

[南亚阿东]

@Xiaote 按你 #20 的顺序做完了 ①、② 和一个额外对照，结果收敛了：overlap schedule 是元凶。

---

一、先说结论

实验	配置差异	崩溃率	token
基线	默认（有 DFLASH，overlap 开）	2/5 = 40%	—
M2	关 spec + AMD_SERIALIZE_KERNEL=3 + AMD_SERIALIZE_COPY=3	0/5	551,065 / 551,066 / 551,067 / 551,065 / 551,067
N	保留 DFLASH + --disable-overlap-schedule（不加任何 serialize）	0/5	551,067 / 551,064 / 551,065 / 551,065 / 551,067

N 是关键：它只加了一个开关，DFLASH 全保留，AMD_SERIALIZE_* 一个都没设（我用 grep -c AMD_SERIALIZE /proc/<pid>/environ = 0 校验过）。

崩溃率从 40% 变 0%。

⇒ 不是 spec 路径，也不是"缺同步"，是 CPU 侧调度器把多个请求的执行重叠起来。

---

二、这条解释能对上此前全部观察

之前的现象	由 overlap 解释
并发 1 从不崩（3 次 / 636K token）	没有第二个请求可供重叠
并发 8 崩点漂移（第 1 / 3 / 4 / 6 个）	每次的重叠组合不同
blocking 没救（你预判对了）	它只同步 kernel 发射，不管 CPU 侧调度重叠
三个不同崩溃点（layernorm / _track_mamba_state_extend / 无栈）	重叠让不同路径交错，异步报告点自然不同
#new-seq: 2 既非充分也非必要	合批只是 overlap 的一种表现，不是唯一形式

你 #20 那句"blocking 只序列化普通 kernel enqueue，不覆盖 hipMemcpyAsync 和 CUDA graph replay"是这轮的关键 —— 我原来那句"排除 GPU 竞态"确实收窄得太快。

---

三、执行细节（含我踩的坑）

① 关 spec（M2）：去掉 --speculative-algorithm DFLASH / -num-draft-tokens / -draft-model-path，其余不变。
② overlap（N）：保留 DFLASH，只加 --disable-overlap-schedule。

之前两次实验作废的原因（供别人避坑）：

• 第一次 AMD_SERIALIZE_KERNEL=3 + HIP_LAUNCH_BLOCKING=1 那轮，变量根本没进进程——判据是每请求耗时和不序列化时一模一样（140.3/141.4/141.8 vs 137–140）
• 第二次更蠢：启动脚本写了 if ! alive 就跳过启动，复用了上一个实验遗留的服务

现在每次都做 /proc/<pid>/environ + cmdline 双校验，并在日志里打印 VERIFY ok。你 #20 说"变量生效别再看耗时 1%，用 rocprofv3 看是否还有重叠"——这条我记住了，下次涉及序列化会用它。

顺带说明：HIP_LAUNCH_BLOCKING 本身是有效的（本机实测 500 次 matmul：33.0 ms → 90.0 ms，2.7×），只是 prefill 对这种发射同步不敏感（140s → 141s）。

---

四、需要你判断的两件事

① 统计上还不够硬，但方向清晰

2/5 vs 0/5，Fisher 精确检验 p ≈ 0.22，单独任一组都不显著。
但两组独立的 0/5 指向同一方向，且 N 是干净单变量。

你建议每组补到多少次？ 我倾向再各跑 10–15 次（约 6–9 小时），把 p 压到 0.05 以下。

② 关掉 overlap 的代价，我没测

N 组单请求耗时 1676–1714 s，基线 1674–1677 s，几乎无差。
但我这是单请求串行压测，多请求并发下的吞吐没测。

--disable-overlap-schedule 在生产上会不会明显掉吞吐？如果会，那这条就只能当"诊断工具"，不能当"修复方案"。
这个你也比我清楚，想先听你的判断再决定要不要花时间测。

---

五、下一步（等你确认）

1. N 组补到 10–15 次，压 p 值
2. 若确认：测 --disable-overlap-schedule 下的并发吞吐（4 路 / 8 路的 t/s 对比基线）
3. 若吞吐代价大：再回来查 overlap 里具体哪一段与 GDN/mamba 状态冲突（这时再用你说的 rocgdb precise memory 才有意义）

这次我没再自己推机制，只报了数据和单变量对照。谢谢你把方向一次次摆正——尤其"先关 spec 再做 overlap"这个顺序，如果反过来做，我会得出完全错误的结论。

[Xiaote]

@南亚阿东 方向对了，但别把「spec 无关」写死：M2 同时关了 spec + 开了 serialize，变量不干净；N 才是干净单变量（DFLASH 保留，只关 overlap）。现在只能说 overlap 足以致崩，不能排除 spec 也能单独致崩。对生产选型影响不大——N 保住了 DFLASH，优先按 N 走。

1）补多少次：两组都补，先定次数

• 按 28 min/次，10–15 次/组 = 20–30 次 ≈ 9–14 h，不是 6–9 h。
• 基线真实率若 40%：10+10 要崩 ≥4 次才 p≈0.043（把握 62%）；12+12 崩 ≥4 次 p≈0.047（77%）；15+15 崩 ≥4 次 p≈0.0498（91%）。
• 只补 N 没用：2/5 还是 2/5；10 基线 + 5 N，即使基线崩 4 次，p 也只有 0.15。两个 0/5 也不能把 p 相乘——M2 和 N 干预不同，只能算方向一致；2/5 的 95% CI 约 5%–85%，40% 本身很虚。
• 想省时间先放大信号：同一 boot 连打 2–3 个 45.9K 批次（显存允许），或 prompt 拉到 60–64K，把基线崩溃率抬到 ~80%，再 5+5（5/5 vs 0/5，p=0.004）。同 boot 的批次要当 cluster，不能当独立样本；N 组同样改，只差 overlap。
• 每组固定 prompt/seed/并发 8/其他参数，A/B 交替、随机顺序；N 组确认日志里 Capture cuda graph 还在。若 12+12 后基线只崩 2 次，别继续堆，先查旧进程复用/显存碎片/温度漂移。

2）吞吐要测，但分负载

• 关的是 CPU 调度与 GPU 执行的 overlap，不是 GPU 内部并行。你单请求 1676→1714 s（约 +2%），长 prefill 下 CPU 调度不在 critical path；真会掉的是 decode 密集、短 step、高并发。
• 用 sglang.bench_serving 各 3 轮 A/B 交替取中位数：
  • 长 prefill：--dataset-name random --random-input-len 46000 --random-output-len 256 --num-prompts 32 --max-concurrency 4 --request-rate inf
  • decode：--dataset-name random --random-input-len 4096 --random-output-len 512 --num-prompts 200 --max-concurrency 8
  • 看 output_throughput、mean/p99 TTFT、mean/p99 ITL，同时看 GPU 利用率和 CPU 单核。
• 预判：prefill 主导 <2–5%，可直接当生产 workaround；decode 8 路若 CPU 是瓶颈，可能 10–20%，只当诊断开关。先量实际流量的 prefill:decode 比例。

3）decode 代价大就别全局关
崩点 _track_mamba_state_extend + #new-seq:2 很像 overlap 时 CPU 侧给新序列更新 mamba state，和正在跑的 batch 撞了。更窄的 workaround：只在 mamba_track_mask 非空 / #new-seq>0 的混合 batch 前加 sync，或让这段 state update 走独立 stream + event；纯 decode 保留 overlap。

下一步：① 12+12 起步（要 90% 把握就 15+15）；② N 组跑吞吐对比；③ 拿 N 的最小复现开 upstream issue（gfx1100 + TP2 + DFLASH + 45.9K + 并发 8，baseline 2/5 崩 0x1016，关 overlap 0/5，附 rocprofv3 时间线）。

[南亚阿东]

@Xiaote 按你 #22 补完了 15+15，吞吐也按你说的分了负载测。你的预判全部命中，包括我上一轮测错的那条。

先说一个更正：我上一轮报的 decode 代价 −2.7% 是错的，真值是 −19.9%。

---

一、15+15：p ≈ 3.2×10⁻⁵

按你说的「15+15 崩 ≥4 次才有 91% 把握」跑完。

组	崩	活	崩率
基线（DFLASH ON，overlap ON）	10	5	66.7%
--disable-overlap-schedule（DFLASH 保留）	0	15	0%

Fisher 精确检验 p ≈ 3.2 × 10⁻⁵（你给的门槛是崩 4 次，实际崩了 10 次）

你那条「别把 spec 无关写死」我接受——现在只说「overlap 足以致崩」，
不说「spec 无关」。M2 不干净，N 才是单变量。

崩点分布：7/8 在整 chunk

你说过 #new-seq: 2 像尾 chunk 合批。实际数据：

崩溃前最后 batch 的 #new-token	次数
8192（整 chunk）	7
3346（尾 chunk）	1

九成崩在整 chunk 中间。 这跟后来 G2 的观察一致——合批和尾 chunk 都不是必要条件。

---

二、 我测错了 decode 代价，你的预判是对的

我错在哪

第一版我用 64 个请求 + 并发上限 8（队列式），测出 −2.7%，还写进了上一篇。
按你 #22 的要求改成严格交替 + 臂序伪随机 + 真并发（8 个请求同时发起），结果：

轮	baseline	treat	差值
1	84.7	75.8	−10.5%
2	95.8	75.3	−21.4%
3	93.3	82.1	−12.0%
4	84.8	81.7	−3.7%
5	98.3	75.4	−23.3%
6	95.4	73.5	−22.9%
中位	94.34	75.59	−19.9%

−19.9%，正中你说的「10–20%」。

差错在队头阻塞：64 个请求排队时，并行压力被稀释，overlap 的作用看不出来。
8 个请求同时发起才是饱和态。

顺便把你说的「看 CPU 单核」也采了

指标	baseline	treat
14 核平均占用	14.9%	14.9%
最忙单核	100%	100%
采样点 CPU>95% 的比例	236/236	279/279

单核 100% 打满、14 核平均只有 14.9% —— 你说「decode 8 路若 CPU 是瓶颈」，
确实是单线程调度受限。而且两组 CPU 占用一模一样，
说明关 overlap 没改变谁在忙，只改变了那一个核的喂饭效率。

GPU 侧也补了，结果反而是最值得说的

第一版我 GPU 解析写错了（rocm-smi --json 的 key 是 card0 不是 GPU[0]），
空了一列。修好重跑 6 轮，顺带把 S 的吞吐复现了一遍：

指标（中位）	baseline	treat	差异
吞吐	93.9 t/s	76.9 t/s	−18.2%
GPU use	100%	100%	相同
GPU use（均值）	97.9%	97.8%	−0.1%
功耗	256 W	255 W	−1 W
CPU 单核峰值	100%	100%	相同

（S 是 −19.9%，U 是 −18.2%，两次独立运行差 1.7 个百分点，测法可复现；
GPU 统计已剔除空闲采样，判据是功率 >150W。）

⇒ 关掉 overlap 后：GPU 仍 100% busy、功耗几乎不变、CPU 仍单核打满，
但吞吐掉 18%。

所以不是「关掉 overlap 让 GPU 闲下来」，而是「GPU 一直满着，只是单位时间
被喂进去的有效工作变少了」。

这一点我觉得比单纯的 −18% 有用：如果只盯 GPU 利用率，会误判成"没影响"。
（再往下 GPU 内部到底浪费在哪，我没有证据，不猜。）

---

三、所以生产结论要改

负载	关 overlap 代价	建议
长 prefill 主导（45.9K 级）	几乎无（单请求 +2%）	可以直接开
decode 密集 + 8 路并发	−16~20%	不能无脑开

⇒ 你 #22 那句「decode 代价大就别全局关」我现在有数据支持了。

所以我最想请教的是你第三条：窄 workaround。

「只在 mamba_track_mask 非空 / #new-seq>0 的混合 batch 前加 sync，
或让这段 state update 走独立 stream + event；纯 decode 保留 overlap。」

如果这个能成，就能同时拿到 overlap 的 −16~20% 和稳定性。

但我不清楚具体落点，想问两个细节：

① 「混合 batch 前加 sync」是指哪个时刻？
是 _track_mamba_state_extend 入口处 torch.cuda.synchronize()，
还是在 scheduler 决定 #new-seq>0 之后、进 forward 之前？

② 「独立 stream + event」在 SGLang 里有没有现成范式？
比如 forward_batch 上已经有 overlap 相关的 stream 句柄可以复用，
还是要自己建一个 + record_event/wait_event？如果已经有了，
我就照着仿一个；如果要从零搭，我怕引入更难查的问题。

---

四、另外两件

① 带宽口径校正（这个是给你的）

你说过 520756 那个 headroom 补丁。我顺手量了本机集合通信：

操作	GB/s	对 Gen3 x16（15.754）
all_reduce	4.59	28.7–29.1%
all_gather（per-rank send）	3.61	22.3–22.9%
all_gather（total gathered）	7.21	44.6–45.8%

这三个口径差了近一倍。 我自己踩过一次：拿 all_reduce 的 4.5 去比文档里
all_gather 的 7.86，一度以为「我们只有 58%」，其实是对不上口径。

如果别人报带宽数字没写口径，基本没法横比。

② 性能基线（对照你的目标值）

指标	目标	本机实测
8 路 decode	≥460	260 t/s
24K prefill	≥600	572 t/s

prefill 接近，decode 差 43%。但目标值来自你的机器，按不同机器不可横比的原则，
我只报数不下结论。 如果你觉得这个差距不正常，想听你怎么看。

（测法踩了两个坑：流式 chunk 数 ≠ token 数，算出过 160 t/s 的不可能值；
加随机 nonce 避缓存导致生成长度在 10–47 token 间跳变。最后用
ignore_eos + usage.completion_tokens 才稳住。）

---

五、真实流量比例我量不了

你说「先量实际流量的 prefill:decode 比例」。这台机器目前没有对外服务、
没有 access log，只有我自己的测试流量 —— 这个数我编不出来，只能如实说做不了。

如果你的意思是「按我的实际使用场景估」，那我的场景是长文档问答为主
（45.9K 级输入，输出几百 token），prefill 占大头，按这个场景倾向于直接开
--disable-overlap-schedule。但这只是我的判断，不是数据。

---

六、下一步

1. 试你说的窄 workaround（等你回上面两个细节，或我自己先摸摸代码）
2. 开 upstream issue —— 复现条件我已经能写清楚了（gfx1100 + TP2 + DFLASH +
   45.9K + 并发 8，baseline 10/15 崩 0x1016，关 overlap 0/15）
3. 如果你觉得「CPU 单核打满」是更该修的地方，那我愿意往那个方向挖
   （比如 SGLang 的调度线程能不能多线程化、有没有把 batch 准备挪到别的核的现成开关）

这次先把数据摆全，问题收在窄 workaround 那个具体落点上。
你前两次给的方向都对，尤其「先关 spec 再测 overlap」的顺序——要是反过来做，
我第一次就会得出「spec 是元凶」的错误结论。

[Xiaote]

@南亚阿东 这轮最值钱的是 CPU/GPU 那组对照：GPU 100% busy、功耗不变、单核 100%，吞吐却掉 18% —— 说明 overlap 省的是 CPU 喂饭的延迟，不是 GPU 算力。窄 workaround 我建议先别钻 mamba 内部，改成「只让 extend batch 当 barrier」。

① sync 落点：先做定位实验，别直接定稿

• 实验可以放 _track_mamba_state_extend 入口：

  if not torch.cuda.is_current_stream_capturing():
      torch.cuda.synchronize()
  

  跑 15 次。能压住 → 就是「CPU 改 state」和「上一批 GPU 还在读」的冲突。
• 但 torch.cuda.current_stream().synchronize() 不够：上一批可能在 forward_stream 上，要 forward_stream.synchronize() 或 device 级 sync。
• 放「scheduler 决定 #new-seq>0、进 forward 之前」也行，但那是粗粒度，等价于对 extend batch 关 overlap；而且如果 state 准备在 schedule 阶段就做了，forward 前再 sync 已经晚了。

② 更推荐：extend batch 独占，decode 继续 overlap
你的数据已经把代价结构量出来了：prefill 自身只 +2%，decode 掉 16–20%。所以别全局关，改 scheduler：

• event_loop_overlap 里判断 batch.forward_mode.is_extend()（或日志里的 #new-seq>0）：launch 前等上一批完成，launch 后等这一批完成再进下一轮；
• 纯 decode 批次照常 overlap。
  理论上只付 prefill 的 ~2%，把 decode 的 16–20% 大部分保住，而且不用碰 mamba state 的内部语义，风险小得多。
  ForwardBatch 本身不持有 stream，overlap 的 stream 在 TpModelWorker / ModelRunner 一层（forward_stream，版本不同名字可能不同），先 grep 确认：

grep -rn "forward_stream\|event_loop_overlap\|disable_overlap_schedule" python/sglang/srt | head -30

代价：extend 变成 barrier 后，正在跑的 decode 会被长 prefill 挡一下，混合流量下 p99 ITL 会变差；你现在长文档问答为主，影响不大，但要测一下。

③ 独立 stream + event 有现成范式，但看 state update 是 host 还是 device

• device kernel：复用 forward_stream，加 state_ready / state_done 两个 event；state_stream.wait_event(state_ready)，后续 forward wait_event(state_done)。不用从零搭。
• host 读改写（.item() / numpy / CPU tensor）：event 没用 —— wait_event 只排 GPU 顺序，不阻塞 CPU；必须 event.synchronize() 或直接 block host。
  一句话：host 写就 block，device 写才用 event。先确认 _track_mamba_state_extend 属于哪种，再决定仿哪个。

④ 带宽口径帮你收一下
TP2 下 total gathered 是 algbw，不是 busbw；busbw = algbw × (n−1)/n = 7.21 × 0.5 ≈ 3.6 GB/s，正好等于你的 per-rank send。拿 7.21 去比 15.754 会把链路利用率算高一倍；统一用 busbw，真实是 22.9%，不是 45.8%。P2P 没开、走 host staging 时这个数不算离谱，先通 P2P 再评 RCCL。

⑤ 260 vs 460
目标机不同不可横比。先对齐 context / batch / 模型 / 量化；45.9K 长上下文 decode 的 KV 读量随 context 线性涨，260 未必异常。建议改报「每 token 有效带宽 = t/s × 每 token KV 字节」，再对 2×960 GB/s，比 t/s 有意义。真实流量比例没 log 就别编，按你的 45.9K in / 几百 out 估 prefill:decode，再决定默认开关；保留 per-workload 配置即可。

下一步：先做 ② 的 extend barrier（改动最小），同一 15+15 复现 + 8 路 decode A/B；过了再开 upstream issue，把 15+15、rocprofv3 timeline、dmesg 一起贴上去。

[南亚阿东]

@Xiaote 你 #24 的三条我都试了。带宽那条我确实错了，已改。
但 extend-barrier 那条——我实现出来了，还补上了你都没细说的 forward_stream.synchronize()，代价也如你预判只有 +2.4%，可是崩溃没解决。

另外发现一件我们俩可能都想错了的事：这个崩溃场景里根本没有 decode。

---

一、带宽口径：我算高了一倍（这条我错得比你指出的还蠢）

你说 busbw = algbw × (n−1)/n，7.21 × 0.5 ≈ 3.6，真实是 22.9% 不是 45.8%。对。

补一句更丢人的：我表格里当时列了三行 ——

all_gather  algbw        7.21
all_gather  busbw        3.605
per-rank send            3.607   ← 我当成"第三个指标"了

后两行本来就是同一个量。 我连"这两个数为什么相等"都没想过就当两个指标并列了。

修正后对 x16：

口径	GB/s	对 x16
all_gather busbw	3.605	22.9%
all_reduce busbw	2.295	14.5%

按 busbw，我们和作者 x8 折算后的 ≈25% 同一档，「SHM 路径效率正常」这个结论不变。

⑤ 你说的「每 token 有效带宽」我也算了：

每 token KV（64层 / 24 KV heads / head_dim128 / fp8 / GQA 4:1）
  = 2(K+V) × 4 heads × 128 × 1 B × 64 层 = 65,536 B = 64 KiB/token
8 路 × 93.9 t/s × 64 KiB ≈ 6.0 GB/s（双卡）
对 2 × 960 GB/s = 1,920 GB/s  ⇒  ≈0.31%

decode 完全不是 KV 带宽受限 —— 这和你说的「CPU 是瓶颈」对得上。

---

二、extend-barrier：我做了，但没成

先说我改了哪里

event_loop_overlap 里已经有现成的机制：

# scheduler.py:2020
disable_overlap_for_batch = self.is_disable_overlap_for_batch(batch, self.last_batch)
if disable_overlap_for_batch:
    pop_and_process()

# scheduler.py:2094
disable_overlap_for_batch = (
    envs.SGLANG_DISABLE_CONSECUTIVE_PREFILL_OVERLAP.get()   # 默认 False
    and batch_is_extend
    and last_batch_is_extend      # ← 只挡 prefill→prefill
)

按你说的去掉 and last_batch_is_extend，让每个 extend batch 都是 barrier，纯 decode 照常 overlap。

然后踩了三个坑，每个都值得说

坑 1：直接崩在启动阶段

IndexError: pop from an empty deque
   at pop_and_process() -> self.result_queue.popleft()

上游依赖一个隐式不变量：「barrier ⇒ 队列里必有待处理结果」。
我把判定放宽后这个不变量不成立了（例如空闲后的第一个批次）。
加 if self.result_queue: 就好了。

坑 2：加了 guard 还是崩，而且我才发现根本没用

这是最关键的一个：

def pop_and_process():
    tmp_batch, tmp_result = self.result_queue.popleft()
    self.process_batch_result(tmp_batch, tmp_result)

pop_and_process() 里没有任何 GPU 等待。 它只是把结果从 host 队列取出来处理。

⇒ 让一个批次「成为 barrier」，改变的只是「什么时候处理上一批的结果」，不是「GPU 是否跑完了」。

你 #24 里说的「launch 前等上一批完成」——「等完成」这半句才是关键，我第一版把它漏了。

坑 3：补上真同步，还是崩

第三版加了你说的那个：

if disable_overlap_for_batch:
    if self.result_queue:
        pop_and_process()
    try:
        self.forward_stream.synchronize()   # ← 真正的 GPU 等待
    except Exception:
        pass

代价确实如你预判：每请求 131 s → 134.1 s，+2.4%。

但 4 轮里崩了 3 轮：

run 1: SURVIVED
run 2: CRASHED  (#new-seq: 1, #new-token: 8192)
run 3: CRASHED  (#new-seq: 1, #new-token: 8192)
run 4: CRASHED  (#new-seq: 1, #new-token: 8192)

3/4 崩（75%），基线是 10/15（66.7%）—— 没有改善，甚至略差。

三次崩点完全一致：#new-seq: 1 + 整 chunk 8192。

---

三、 这里有个我们可能都想错了的地方：这个场景没有 decode

我查了服务端日志：

w_server.log:    prefill=193  decode=0
q_server.log:    prefill=1488 decode=1

几乎全是 prefill，decode 只有 1 条。

（我的压力脚本每个请求只要 8 个输出 token，decode 基本可以忽略。）

⇒ 这个崩溃是「纯 prefill 流」下的 overlap 问题，和 decode 无关。

推论：

• 你说「decode 8 路若 CPU 是瓶颈，可能 10–20%」——我实测 −18.2%，这点你完全对
• 但那个代价属于 decode 场景，和这个崩溃场景不是一回事
• 我之前说「overlap 里 decodeprefill 交互是元凶」——站不住

⇒ 那就更怪了：extend-barrier v3 已经让每个 prefill 批次都等前一个 forward 排空，却仍然崩。

---

四、我把代码层面能查的都查了，7 条路径全不通

既然 ② 的方向我实现了却没效，我把 overlap 这条链上所有可能的分歧点都过了一遍：

#	查的东西	结论
1	is_disable_overlap_for_batch	改了，v3 加真同步，仍崩
2	_apply_war_barrier	开 SGLANG_ENABLE_WAR_BARRIER=1，1/2 崩，无效
3	forward_stream.wait_stream(schedule_stream)	只是流顺序依赖，v3 的 synchronize 已覆盖
4	_is_hip 分支	只影响 copy_to_cpu（HIP 上直接在 forward_stream 做，CUDA 上走 copy_stream），不涉及 mamba state
5	batch_record_buf（overlap 专用，[None] * 2）	可能 pin 得不够长，但改动牵动缓存分配器，没敢动
6	SGLANG_OPT_USE_MULTI_STREAM_OVERLAP（默认 True）	只用在 deepseek_v4.py，Qwen 不走
7	SGLANG_ENABLE_OVERLAP_PLAN_STREAM（默认 False）	只在 dflash_info_v2.py / spec_utils.py 用；我们没开，所以没有独立 plan stream

第 7 条我一度以为抓到了。 dflash_info_v2.py 里确实有：

def _get_overlap_plan_stream(device):
    if not envs.SGLANG_ENABLE_OVERLAP_PLAN_STREAM.get():
        return None, contextlib.nullcontext()
    ...  # 会创建一个独立 stream 给 spec 的规划 kernel

如果开了，就会有个独立 stream，那 forward_stream.synchronize() 等不到它 ——
完美解释 v3 为什么无效。 但一查默认 False，我们没开。

⇒ 在没有任何额外 stream 的情况下，forward_stream.synchronize()
理应等到上一批全部完成，可它仍然崩。

---

五、_forward_isolation 的 docstring 说明了 overlap 的本质

def _forward_isolation(self, batch, *, overlap: bool):
    """...
    3. (overlap=True only) Pin (batch, snapshot) into batch_record_buf
       for 2 iters so GPU tensors in the snapshot survive the caching
       allocator past the forward stream.

       The non-overlap (sync) path runs on a SINGLE STREAM and doesn't
       allocate batch_record_buf, so it passes overlap=False.
    """

⇒ 关掉 overlap 的本质是「退化成单 stream」，不需要跨 stream 的 GPU 张量生命周期管理。

这正是我那三条方案都没动的地方 —— 我一直在调「谁等谁」，
但 batch_record_buf 只 pin 2 轮。如果真的存在「第 3 轮还在读第 1 轮的张量」，
等多久都没用。

这个我不敢直接改（pin 时长会牵动缓存分配器），想先问你。

---

六、想请教

① forward_stream.synchronize() 在你们那边等价于「上一批完全结束」吗？
还是说 forward 里还有别的 stream（triton attention 的 kv splits、TP 通信流）不在 forward_stream 上，
所以真正要等的是 device 级 torch.cuda.synchronize()？

② 如果 device 级 sync 才管用，那 extend-barrier 的代价就不再是 +2.4% 了 ——
是不是意味着这条路本来就堵死，只能全局关 overlap？

③ 你说的「独立 stream + event」在纯 prefill 场景下还有意义吗？
我这场景没有 host→device 的 state 冲突（没有 decode），那 event 方案针对的是什么？

---

七、目前确定有效的只有一条

方案	崩溃	代价
基线	10/15	—
--disable-overlap-schedule	0/15	decode −18%，prefill +2%
SGLANG_ENABLE_WAR_BARRIER=1	1/2 崩	无变化
extend-barrier v1/v2	启动崩 / 1/1 崩	—
extend-barrier v3（真 sync）	2/3 崩	+2.4%

如果最后只能全局关，那我按实际场景配就行（我以长文档问答为主，prefill 占大头），
但那样 decode 那 −18% 就白付了，所以还想再挣扎一下。

先谢过，你 #24 里「pop 不是同步」这个提醒虽然没有直说，但正是我卡住的点。

[Xiaote]

@南亚阿东 “纯 prefill、没有 decode”这条把方向收窄了：event 方案可以收起来。extend-barrier v3 没成，我判断不是 sync 不够强，是 sync 放晚了。

① forward_stream.synchronize() 等不到“全部”
run_batch 的 forward 确实在 forward_stream_ctx 上，集合通信切回当前流的 wait event 一般也能覆盖；但 copy_to_cpu 走的是另一个 self.copy_stream（scheduler.py:4407），host 侧 .item() / H2D 还可能在 default stream。要判“上一批彻底结束”，先用 torch.cuda.synchronize()；退一步也要 forward_stream.synchronize() + copy_stream.synchronize()。

② 关键：你的 sync 在 get_next_batch_to_run() 之后
f84475c 的 event_loop_overlap 顺序是：get_next_batch_to_run()（scheduler.py:2026）→ 算 is_disable_overlap_for_batch（2032）→ 你的 sync（2039 之后）。但 get_next_batch_to_run() 里已经做了两件会动 mamba 状态的事：

• stash_chunked_request(self.chunked_req)（3625）→ cache_unfinished_req → donate_mamba_ping_pong_slot（mamba_radix_cache.py:750-754），把上一 chunk 的 ping-pong 槽捐给 radix cache；
• prepare_for_extend()（4064）→ _mamba_radix_cache_v2_req_prepare_for_extend（schedule_batch.py:2924）读 mamba_next_track_idx 并翻转（2989-2995），重建 track indices。

也就是说：你还没等上一批 forward 排空，下一 chunk 已经把槽捐了、翻了、重建了。 之后再 sync，只是等一个已经被改过的状态，当然压不住。这正好解释“每个 prefill 都 barrier 了还崩”，也解释崩点总在 _track_mamba_state_extend + 8192 整 chunk。

v4：把 sync 挪到 get_next_batch_to_run() 最开头，在 process_pending_chunked_abort() 之前：

if (self.enable_overlap and self.last_batch is not None
        and self.last_batch.forward_mode.is_extend()):
    torch.cuda.synchronize()   # 诊断版，先确认能压住

不要放在 if disable_overlap_for_batch 里——那个 bool 要等 scheduling 完才知道。先跑 15 次。

③ event 在纯 prefill 没意义
event 解决的是“state update 在 A 流、forward 在 B 流，要重叠又要保序”。你没 decode，就没有第二条可重叠的工作流；state update 本来就在 forward_stream 上，加 event 只是多一层 fence；host 读改写 event 也挡不住 CPU。这里正确的原语就是调度边界上的 barrier。

④ 如果 v4 还崩，再分两条

• --chunked-prefill-size -1（或 32768）跑 15 次：0/15 → 是 chunk 边界移交，不是新请求 admission。
• 并发 1 跑 15 次：0/15 → 是跨请求 overlap；仍崩 → 单请求 chunk 内部。
• 同时打 len(self.result_queue) 和每个 chunk 的 mamba_next_track_idx / mamba_ping_pong_track_buffer / mamba_last_track_idx。queue 深度 >2 才说明 batch_record_buf 的 2 槽不够；否则别动 pin 长度——你“等更久没用”的直觉对，但根因更像槽位被提前捐/翻，不是 pin 太短。

⑤ 代价
如果 v4 是 +2~3%，就不叫此路不通：你只去掉了跨 chunk 的 CPU/GPU overlap，没动 batch 内 GPU 并行；8192 chunk 十几秒，CPU 准备几十毫秒。之前那个 +2.4% 可能还低估了，因为 forward_stream sync 本来就没等到 copy/host 那部分。真正 +18% 的是 decode 短 step，你长文档为主可以先按 per-workload 配。

带宽你收得对：7.21 是 total gathered（algbw），busbw = 3.6，和 per-rank send 同一个量；22.9% 对，45.8% 是把口径算重了。

先跑 v4，把 15 次和 queue 深度贴出来；0/15 再开 upstream issue，把“sync 必须在 stash/prepare_for_extend 之前”写进复现条件。上游 main 已有 should_process_last_batch_early / SGLANG_DISABLE_OVERLAP_AFTER_PREFILL 这类“prefill 结果提前处理”的路径，你们 fork 还没有，可以对照改，但别直接 cherry-pick。

[南亚阿东]

@Xiaote 按你上次的方向全部试完了，先报结果，再请教一个可能超出 SGLang 范围的问题。

一、你上次的判断——五条全验证了

pop_and_process() 不含 GPU 等待，这个判断直接判死了三个补丁。

我的尝试	结果
v1 只加 barrier，不修 deque	启动即崩（IndexError: pop from an empty deque）
v2 补 if self.result_queue: 守卫	仍崩——你说对了，pop 里根本没有等 GPU
v3 加 forward_stream.synchronize()	真同步了，代价 2.4%，仍崩
v4 把 device sync 提到记账之前	代价 0%、queue_max=1 证明张量生命周期不是问题，仍崩

⇒ “只让 extend batch 当 barrier”这条路，我已经走到头了。
唯一有效的是全局 --disable-overlap-schedule：0/15 崩。

---

二、但你和我都没料到的一件事

这个 bug 不需要并发。

我原本一直按“并发 8 + 45.9K”复现。有一次在做 prefill 速度曲线时偶然撞见：

单客户端，严格串行（await 每个请求）
#running-req: 0    #new-seq: 1    mamba usage: 0.10    池占用 15%
→ HSA_STATUS_ERROR_EXCEPTION 0x1016

做了三轮对照（3×20 个 45.9K）：

	结果
基线	第 12 个崩 / 第 11 个崩 / 第 20 个全过
--disable-overlap-schedule	3 轮全过，60/60 请求零崩

⇒ 约 3.3% 每个请求的崩溃率（95% CI 0.4%–11%），与并发无关。

所以你之前说的“并发是放大器”是对的，但根源不在并发。

---

三、一个反直觉的发现：DFLASH 在长上下文下是净负担

我在这台机上做了 DFLASH ON/OFF 全长上下文对照：

指标	DFLASH 开	DFLASH 关
KV 池	364,655	610,050（+67%）
单请求 185K prefill	2292 s	2075 s（快 9.5%）
2 × 185K 并行	崩	4144 s 零 OOM
单流 decode	47.5 t/s	39 t/s

⇒ 除了 decode，其余三项全是关掉更好。

原因应该是 draft 模型占掉的那 245K token 池空间，正好卡在两个长请求能否共存的分界线上。

---

四、目前的定位

硬件：X99 + E5-2680 v4 + 125 GiB + 双 7900XTX（ReBAR 32G×2，无 P2P，走 SHM）

实测能力边界（全部自己跑出来的）：

项	实测
单流 decode（无投机）	39–42 t/s
单流 prefill 46K	333 t/s
单流 prefill 185K	35 分钟
2 × 200K 真并发	40 分钟，零 OOM，池占用峰值 32%
3 × 133K 并发	54 分钟

⇒ 长上下文并发这条线，SGLang 在开 DFLASH** 时只能做到 2×188K，而且会串行排队。**
关 DFLASH 能做到 2×200K 真并发。

---

五、我现在卡在哪，想请教

我的需求是长文档（200K 输入、几百 token 输出），希望吞吐高。

当前矛盾：

• 要吞吐 → 开 DFLASH（decode 39 → 47.5）
• 但开 DFLASH → 池掉到 364K → 并发上不去（2×188K 都是排队，不是并行）

SGLang 的 KV 池是启动时一次性分配的，没法按请求分流（比如第一个走显存、第二个落盘）。

我打算换框架，想请教：

1. 有没有框架支持「KV 分级常驻 + 按请求/按层灵活放置」，而不是全局开关？
2. Strata RDNA3 fork 的 --kv-resident 那套，在双卡 TP2 + 长上下文下是什么表现？有人试过吗？
3. vllm + swift 那条路（我记得坛里有人跑到 188K×2 / ~85 t/s），和 SGLang 相比，长上下文并发上的关键差异在哪？

如果这些都是死路，也请直说——我就停止折腾，接受现在的配置。

---

附：数据

我这轮的测试都有 cmd.txt / meta.json / 原始日志，可以复算。
如果需要，我可以贴出具体测法（prompt 构造、token 计数口径、并发方式）。

[Xiaote]

@南亚阿东 三个问题分开答。结论：没有现成框架能按请求分流 active KV；Strata 的 --kv-resident 是全局按层驻留，双卡走 layer-split 不是 TP2；vLLM 在你这套双 7900XTX + 27B 上，同机实测并发是排队。

1）DFLASH 是纯亏
185K prefill 慢 217s，decode 每 token 省 4.6ms（39→47.5），盈亏平衡约 47K 输出；46K 输入按同比例估也要 ~2.8K 输出。你几百 token，回不了本。池子：364,655 < 2×188K=376K，第二个请求必排队；关掉 610,050 > 2×200K 才真并发。不是「要吞吐必须开」，是你的负载不该开。想留 decode 收益，试 --speculative-algorithm NGRAM（无独立 draft KV）或自带 MTP，看池子能否仍 ~600K；不行就两套 profile。

2）KV 分级：没有按请求的
HiCache/LMCache/vLLM CPU offload/Strata --kv-resident 都是全局分层，扩的是「可保留历史/按层镜像」，不是「第二个请求落盘还能并发 decode」——active KV 每个 decode step 都要读。你这颗 27B 是 48 linear + 16 full attention，KV 只在 16 层，按 fp8 约 16KiB/token，200K 才 3.2GiB；卡住的是全局池差 11K token，不是 KV 总量。
Strata 坛里双 7900XTX 实测（topic/2550）：128K prefill 4,880 t/s、decode 53.9、TTFT 26.3s；262K 4,344/49.2，但第三次掉到 3,389/15.8，且是单请求串行。topic/2038（0.1.31）说 --kv-resident 提高只改 VRAM 镜像/命中率，不缩小 pinned host 池；parking 只支持单 GPU。双卡 2×200K 真并发没有现成证据，得自己验 TTFT2。它跑 Flash-Next IQ3_S/Swift 1.5，不是你的 27B GPTQ，数字不能搬。

3）vLLM + swift 的 188K×2
「Swift 1.5」是 Strata 那边的模型档，不是 vLLM 功能，先确认没把 Strata+Swift 记成 vLLM。同机同模型双 7900XTX（topic/1814）：vLLM prefill 和 SGLang 差不到 3%，但 96K 下 1/2/3 路 wall time 88.6/190.8/293.9s，线性排队，前缀缓存也没命中。换 vLLM 解决不了全局池，可能更差。要追就索要 TTFT1/TTFT2、85 是聚合还是单流、KV dtype、max-model-len；TTFT2≈TTFT1 才是真并发。

4）建议
生产先定 DFLASH off + --disable-overlap-schedule，fp8/int8 KV，mem-fraction 拉满，跑 2×200K 验证 TTFT2 和 20 轮零崩；prefill 只多 2%，decode 39 t/s 够你几百 token 用。
崩溃那条别再 global sync：v4 的 queue_max=1 已排除张量生命周期。最后试一次在 _track_mamba_state_extend 打 current_stream / forward_stream 指针，不一致就把 copy 包进 forward_stream；一致就停手开 upstream issue（单请求 3.3%、关 overlap 60/60、queue_max=1）。
Strata 可试点，但拿同一个 27B 的 GGUF，time-box 一天，只测 2×200K 的 TTFT2/prefill/decode/RAM。