Qwen3 C++ Reference Runtime

本目录保存用于提炼 MetaInfer Contracts、Notebooks、Golden Data 和 Immutable Oracles 的可执行 Reference Implementation。Generator Agent 不得读取、复制或链接本目录源码。

可审计性能证据

公开仓库只保留能够支撑主要性能结论的精选原始结果，而不提交模型权重、编译目录、完整 profiler trace 或临时日志。测试口径、原始轮次、脱敏规则和 SHA256 清单见 benchmarks/。这些结果用于同一实现不同 checkpoint/路径之间的 A/B 分析，不构成与 vLLM、TensorRT-LLM 等外部框架的横向基准。

M0 前置条件

CMake 3.28+
安装在 /opt/dtk 的 Clang 15 和 DTK/HIP 4.4
通过 render Group 访问 Z200SM/gfx906 Device
仅供 Offline Tooling 使用的 Python 3.9+
用于 No-network Build Gate 的 strace

依赖同步

只有以下显式 Sync 命令允许访问网络：

python3 tools/vendor_dependencies.py sync

同步后使用离线检查确认 Vendor Tree、License 和内嵌依赖没有漂移：

python3 tools/vendor_dependencies.py check

正常 Configure、Build、Test 和 Runtime Path 必须完全离线。

M0 验收

bash tools/verify_m0.sh --host-only
bash tools/verify_m0.sh

--host-only 适用于当前无法访问 GPU 的环境；仍要求 /opt/dtk 中的 Clang 15。完整验收要求 Host/HIP CTest 通过、两套 Build 都没有建立 AF_INET/AF_INET6 Socket、真实 gfx906 Hardware Probe 通过，并且 C++ Generator Source-isolation Test 通过。

M1 Core 与 Backend

M1 增加以下原生运行时契约：

Status / Result<T> 显式错误传播；
经 Shape、Stride、Alignment 和 Byte Range 校验的 TensorView；
move-only Buffer Host/Device Storage Ownership；
确定性 CPU Allocation、Copy、Stream 和 F32 row-major GEMM；
原生 HIP Allocation、H2D/D2H Copy、non-blocking Stream 和 hipBLAS F32 GEMM。

M1 的 BLAS Gate 有意只开放 F32。BF16 必须在 M3 通过当前 DTK/hipBLAS 的真实 Capability 和 Numerical Test 后才能进入 Model Execution，当前路径对 BF16 返回 Unsupported，不会回退到 CPU。

bash tools/verify_m1.sh --host-only
bash tools/verify_m1.sh

完整 M1 Gate 会先执行 M0 的离线构建、网络 Trace、Source Isolation 和 Device Probe，再在真实 gfx906 上运行 Allocation、Copy、Stream 和 hipBLAS Test。

M2 Model Inputs 与 Weight Loading

M2 固定支持官方 Qwen/Qwen3-0.6B 的 immutable revision c1899de289a04d12100db370d81485cdf75e47ca。模型文件只保存在 ${HOME}/.cache/metainfer/models/Qwen3-0.6B，不进入 Git。

python3 tools/sync_model.py sync --endpoint https://hf-mirror.com
python3 tools/sync_model.py check

Runtime 严格校验 Config、固定 Chat Format、官方 Tokenizer 逐 token fixture，以及 311 个 BF16 Safetensors 的 name、shape、dtype、offset、bounds 和 alias。所有验证通过后，权重以 256 字节对齐布局进行一次后端分配和一次最终同步；失败时不发布部分结果。

bash tools/verify_m2.sh --host-only
bash tools/verify_m2.sh

完整 M2 Gate 会在真实 Z200SM/gfx906 上将 1,503,264,768 字节权重上传到单个 HIP allocation，并继续执行 M0/M1 的 Offline、Host、HIP 和仓库回归检查。

M3 Golden Oracle（仅生成时）

M3 的日常 Build、CTest 和 Runtime 不依赖 Python 推理框架。官方 Golden 只在独立 cache 中用固定的 CPU oracle 生成一次，提交后由标准库 validator 离线检查：

.venv/bin/python -m pip install \
  --target ~/.cache/metainfer/oracle-envs/qwen3-m3-packages \
  --index-url https://download.pytorch.org/whl/cpu torch==2.7.1+cpu
.venv/bin/python -m pip install \
  --target ~/.cache/metainfer/oracle-envs/qwen3-m3-packages --upgrade \
  transformers==4.53.2 safetensors==0.5.3 numpy==2.5.1

PYTHONPATH="$HOME/.cache/metainfer/oracle-envs/qwen3-m3-packages:$PWD/tools" \
  .venv/bin/python tools/generate_m3_golden.py \
    --model-dir "$HOME/.cache/metainfer/models/Qwen3-0.6B" \
    --output-dir testdata/m3
python3 tools/check_m3_golden.py \
  --metadata testdata/m3/official_qwen3_0_6b.json

Artifact 只包含固定 token/position ID、八个小型 F32 capture、Top Logits 和四个 Greedy Token ID；不包含 Prompt、生成文本或完整 Weight。生成环境与包 cache 均不进入 Git，也不进入 Runtime 的 Allow-list。

M3 Numerical Execution 验收

bash tools/verify_m3.sh --host-only
bash tools/verify_m3.sh

M3 的 CPU/HIP 算子在官方 eager graph 产生 BF16 tensor 的边界执行显式舍入， FP32 reduction、softmax statistics 和 GEMM accumulation 保持 FP32。完整 Gate 在真实 Z200SM/gfx906 上验证原生 HIP operators、tiny graph、八个官方 capture、Top Logits 及四个 Greedy Token，并审计新增 Host/HIP 路径没有建立 AF_INET/AF_INET6 socket。首层 capture 使用严格 BF16 activation tolerance；final_hidden 使用独立的深层累积 tolerance，同时继续由 logits、top-32 overlap 和 exact greedy sequence 约束最终行为。

M4a Paged KV 与动态批处理

M4a 将 Qwen3 的增量执行状态移入预分配的 PagedKvCache。Cache 使用一次 Backend allocation 保存 BF16 K/V pool；K 与 V 分区均按 256 字节对齐，逻辑布局为 [layer, physical_block, token, kv_head, head_dim]。每个物理 block 只属于一个 sequence，BlockHandle 携带 generation，旧 view、重复释放和跨 stream 释放都会被拒绝。

KV 容量变更遵循显式事务：Reserve 要么取得本次需要的全部 blocks，要么完整回滚； Commit 原子发布新 block table；Rollback 归还未发布 blocks。Advance 只在 paged forward 的 logits 已复制并完成 owning-stream 同步后推进 token count。终态释放同样先同步 owning stream，再递增 generation 并回收 blocks。

Scheduler 每个 tick 根据当前请求状态重建不可变 StepPlan：先为已有 decode 请求保留一个 token budget，再按 FIFO 接纳能够完整放入本 tick 的 unchunked prefill。 InferenceEngine 按 plan 依次执行 prefill 和 decode，通过最低 token ID 打破 greedy 并列，并在 Apply 后发布事件。完成、取消、模型失败和 shutdown 都进入同一终态清理路径；如果 stream 同步暂时失败，Engine 会保留 pending terminal record 并在下一 tick 开始时重试，避免丢失事件或泄漏 KV blocks。

bash tools/verify_m4a.sh --host-only
bash tools/verify_m4a.sh

Host Gate 在 sanitizer build 中覆盖事务回滚、generation、sequence 隔离、4,096 次复用、10,000 tick scheduler stress、动态成员变化和终态清理。完整 Gate 还会在真实 Z200SM/gfx906 上运行原生 paged attention、tiny paged runner、dynamic engine 和官方 Qwen3-0.6B inspector；后者逐步比较 contiguous/paged logits 与四个 greedy token，验证 steady-state allocation 计数不变且所有 KV blocks 被回收。M4a 新增的 Host/HIP 路径均由 strace 审计，不允许建立 AF_INET/AF_INET6 socket。聚合证据写入被 Git 忽略的 build-host/m4a-evidence.json 与 build-hip/m4a-evidence.json。

M4a 有意不包含 partial/chunked prefill、preemption、CPU swap、随机采样、HTTP 服务和 profiling；chunked prefill 在 M4b 中实现，sampling 与服务边界在 M5 中实现。

M4b Chunked Prefill

M4b 允许 prompt 长度超过单 tick token budget。Scheduler 为 request 一次性预留完整 prompt + max_new_tokens KV 容量，并独占 prefill_cursor；每个 immutable plan item 只携带当前 token slice、absolute positions 和 completes_prompt。配置默认 block_size_tokens=16、prefill_chunk_tokens=256，chunk 上限必须是 block size 的正整数倍，Scheduler 与 cache 的 block size 必须一致。

每个 tick 先为所有 active decode 各分配一个 token，剩余 budget 才分给 partial prefill。每个 request 每 tick 最多执行一个 slice；受限 budget 下，未执行和已执行的 request 通过 prefill deque 跨 tick 轮转，避免固定队首长期占用。非最终 slice 按 block 边界截断，最后的 prompt remainder 可以短于一个 block。

Engine 对中间 chunk 只执行 ForwardPaged、有限值检查和 Scheduler Apply，不会做 greedy，也不会发布 token。只有 completes_prompt=true 的最终 chunk 才产生首个 token；之后进入普通 decode。中途取消、模型失败和 shutdown 继续复用 M4a 的 retryable terminal cleanup，只有 owning stream 同步成功后才发布终态并回收 blocks。

bash tools/verify_m4b.sh --host-only
bash tools/verify_m4b.sh

完整 M4b Gate 会递归运行 M0-M4a，并在真实 Z200SM/gfx906 上把固定 32-token fixture 扩展成 288-token 长前缀。Inspector 分别用 128 和 256 chunk 与 contiguous logits/ greedy 对照，还要求短请求 decode 能在长 prefill 未完成时前进、steady-state hipMalloc 计数不变、KV 全量回收且无 AF_INET/AF_INET6 socket。聚合证据写入被 Git 忽略的 build-host/m4b-evidence.json 与 build-hip/m4b-evidence.json。

M4b 不包含 preemption、CPU swap、随机采样、HTTP 服务、profiling 或多 stream；采样与服务边界由 M5 提供。

M5 Sampling 与 OpenAI Service

M5 的 SamplingConfig 支持 temperature、top_k、top_p、 repetition_penalty 和显式 seed。处理顺序固定为 repetition penalty、 temperature、top-k、top-p、categorical sample；同分时 token ID 较小者优先。 temperature=0 是完全绕过 RNG 的 greedy path。随机路径使用 RNG(request_seed, generated_token_index)，不会读取 batch 序号、request ID 或其他请求状态，因此相同 seed 在不同 admission/batch order 下保持相同结果。

Stop string 在 Service Submit 时由同一 Qwen3 tokenizer 转成 token sequence。 Scheduler 暂存仍可能成为 stop prefix 的 token；完整匹配时不发布 stop，prefix 失配或 length terminal 时才释放安全 token。EOS 同样不进入 HTTP 输出。Terminal event 明确区分 stop 和 length，并携带 prompt/completion token usage。

HTTP 层使用 vendored cpp-httplib 和 nlohmann/json，提供：

GET  /health
GET  /metrics
GET  /v1/models
POST /v1/chat/completions
POST /v1/completions

Chat 和 Completion 都支持普通 JSON 与 stream=true SSE；SSE 最后固定发送 data: [DONE]。HTTP worker 只通过 ServiceRuntime 的有界 command/event queue 执行 Submit、Poll、Cancel；单 Engine worker 独占 Scheduler、Runner、Paged KV 和 compute stream。Client disconnect 会触发 cancel，完成、取消、失败和 shutdown 继续走同一条幂等 KV cleanup 路径。

请求 body、connections、queued connections、admission/pending requests、context、 max output 和 per-request event buffer 都有启动上限。非法字段/类型、unknown model、 context overflow 返回 structured error；admission 满返回 429，starting/draining/stopped 返回 503。Runtime 和 HTTP metrics 不记录完整 prompt、generated text、token array 或 weight content。

HIP build 后可启动真实 Qwen3-0.6B 服务：

./serve.sh --host 127.0.0.1 --port 8000 \
  --max-context-tokens 4096 --max-output-tokens 256

示例请求：

curl -sS http://127.0.0.1:8000/v1/completions \
  -H 'Content-Type: application/json' \
  -d '{"model":"Qwen3-0.6B","prompt":"Hello","max_tokens":8,"temperature":0}'

M5 Gate：

bash tools/verify_m5.sh --host-only
bash tools/verify_m5.sh

完整 Gate 在真实 Z200SM/gfx906 上启动 loopback JSON/SSE service，验证四并发请求、 JSON/SSE 等价、batch-order-independent sampling、stop suppression、disconnect cancel、 shutdown 幂等、steady-state hipMalloc 不增长和 KV 全回收。证据写入 Git 忽略的 build-host/m5-evidence.json 与 build-hip/m5-evidence.json。Fault injection、soak、 native profiling、SIGTERM deadline 和 performance baseline 属于 M6。

M6 Reliability 与 Performance Evidence

M6 不扩大模型或 API 功能面，而是在 M5 真实服务路径上补齐 G6/G7。测试专用 FaultInjectingBackend 可在 Allocate、Copy、GEMM 和 Synchronize 边界确定性失败； Request Error 只结束受影响请求，Fatal Error 停止 admission，所有终态继续复用同一条幂等 KV cleanup 路径。ServiceRuntime::Shutdown() 默认等待 10 秒，超时保持 Draining；真实进程的 SIGINT/SIGTERM 顺序固定为 Stop HTTP Admission、Drain/Cancel、 Flush Profiler、Release，deadline 超时使用进程退出码 124 避免无界析构。

原生 Profiler 默认关闭。启用时必须同时设置三个环境变量：

export METAINFER_PROFILE=1
export METAINFER_PROFILE_OUTDIR="$PWD/build-hip/profile"
export METAINFER_PROFILE_DURATION_S=300
./serve.sh --shutdown-timeout-ms 10000

事件 buffer 在启动时一次性预分配，满后只增加 dropped counter。Operator、Engine、 Service 和 HTTP 只记录固定数值字段；HIP Operator span 表示 enqueue，不增加逐算子 synchronize。Graceful shutdown 原子写出 metrics.json 和 Chrome Trace-compatible trace.json，两者禁止出现 prompt、生成文本、token array、weight 或 tensor content。 /metrics 同时公开 TTFT、TPOT、E2E、executed token、KV current/high-water 和 profile recorded/dropped counter。

M6 Gate 有三个档位：

bash tools/verify_m6.sh --host-only  # Fault/Profiler/Shutdown 与 tooling
bash tools/verify_m6.sh --quick      # 真实 GPU，至少 120 秒，非 certification
bash tools/verify_m6.sh              # 真实 GPU，固定 1800 秒 certification

开发阶段可用 METAINFER_REF_M6_SOAK_SECONDS 显式缩短实机 workload，但任何环境覆盖都会写出 certification=false。完整 workload 覆盖短/长 prompt、JSON/SSE、greedy/采样、 disconnect、非法请求、admission saturation，以及活动 SSE 下 SIGTERM。Gate 要求终态 accounting 一致、KV 全回收、warmup 后 allocation call 不增长、零 dropped profile event、真实 gfx906 operator activity 和有限的非零 throughput；不使用跨机器硬性能阈值。

机器相关结果保存在 Git 忽略的 build-hip/m6-performance-baseline.json、profile 目录和 build-hip/m6-evidence.json。完整性能方法记录 concurrency 1/4/16、prompt token 32/256/1024、output token 1/32/128 与 prefill-only/decode-heavy/mixed 矩阵；日常 G7 Gate 运行每类 representative scenario，完整矩阵由独立 baseline 运行扩展。

M6 明确不加入 preemption、CPU KV swap、tensor parallel、量化、fused attention、 speculative decoding 或多 compute stream。M7 只消费 M0-M6 的 executable evidence，导出 generator oracle cases，并据此更新 Prompts、Notebooks 和 contract traceability；未经 Reference Test 验证的结论不得回流生成链路。

M6.5 Performance Hardening

M6.5 将 reliability soak 与性能测量彻底分开。普通单配置 CMake 构建默认使用 Release；verify_m0.sh 显式使用 Debug Host sanitizer build 和 Release HIP build。性能 driver tools/run_m65_benchmark.py 在 profiler-disabled server 上持续提交 SSE 请求，以客户端 token timestamp 计算 TTFT、TPOT、E2E 和 active-window throughput，不包含固定 soak sleep。正式比较固定运行五轮并使用五轮中位数。

M6.5-A 在 commit 5ba5f09 生成 Release before baseline，在 commit ca08633 测量 trusted paged-attention、hipBLAS stream-binding cache，以及 gfx906 单 wavefront paged softmax reduction 修复。两者使用同一模型 revision、device 0、server limits 和场景矩阵：

场景	指标	Before	After	变化
latency-short, C1	TPOT p50	100.073 ms	99.857 ms	-0.22%
latency-prefill, C1	TTFT p50	288.247 ms	320.220 ms	+11.09%
batch-decode, C4	TPOT p50	400.627 ms	399.998 ms	-0.16%
batch-decode, C4	completion throughput	9.964 token/s	9.986 token/s	1.002x
saturation, C16	TPOT p50	1600.589 ms	1597.428 ms	-0.20%

因此 M6.5-A performance Gate 为失败，不能将去同步描述为主要性能突破。完整 Gate 的 Host/HIP/tooling、真实 Qwen3 shape、chunk equivalence 和 M6 quick regression 均通过，但 TTFT、TPOT、throughput、p95 和 direction-consistency 性能条件没有同时满足。实现保留 checked Paged Attention 的非法表检测，并为 PagedKvCache 来源增加 trusted fast path；该修改清理了 layer-local synchronization boundary，但最终 logits D2H 同步仍需等待全部有数据依赖的 GPU 工作。测量证明主要瓶颈是 Engine 逐序列调用 Runner：C1、C4、C16 的 TPOT 近似按 1:4:16 放大，而 aggregate throughput 基本维持 10 token/s。

Git 忽略的原始证据位于 build-perf/m65-before-summary.json、 build-perf/m65-after-summary.json、build-perf/m65-runs/ 和 build-perf/m65a-evidence.json。before/after summary SHA256 分别为 3b1265b33a5f6a4ff740babf6b5b9452987da738c24799255ab2b67973fae824 和 1aaad9024f85d84666a3bde301c3144809b3dd99a9cb8d4914f1f01e825ac650。

M6.5-B 在 commit bdea0cc 将同一 tick 的 decode rows 组装为一个严格验证的 packed batch。Runner 对整个 batch 只执行一次模型图、一次 logits D2H 和一次最终 synchronize， batched paged attention 按 row 隔离 block table；KV 在 batch forward 成功后原子推进， sampling、Scheduler apply 和事件发布仍保持原有稳定顺序。相同 before baseline 与五轮 Release after 测量得到：

场景	指标	Before	M6.5-B	变化
batch-decode, C4	TPOT p50	400.627 ms	101.832 ms	-74.58%
batch-decode, C4	completion throughput	9.964 token/s	31.440 token/s	3.155x
batch-decode, C4	E2E p95	6422.141 ms	2034.955 ms	-68.31%
batch-mixed, C4	TPOT p50	422.142 ms	118.979 ms	-71.82%
batch-mixed, C4	completion throughput	7.730 token/s	15.118 token/s	1.956x
saturation, C16	TPOT p50	1600.589 ms	108.332 ms	-93.23%
saturation, C16	completion throughput	9.956 token/s	51.505 token/s	5.173x

M6.5-B performance Gate 通过：C4 TPOT p50 改善超过 30%，C4 completion throughput 超过 1.5x，两项均为五轮中的 5/5 轮方向一致；全部场景 TTFT/TPOT p95 的最大回退为 latency-prefill TPOT 的 +5.29%，低于 10% 上限。真实 service inspection 记录 decode_batch_calls=15、max_decode_batch_size=4、allocation calls 3 -> 3 且 KV 全回收。

Git 忽略的 M6.5-B 证据位于 build-perf/m65b-after-summary.json、 build-perf/m65b-service-inspection.json、build-perf/m65b-evidence.json 和同一 build-perf/m65-runs/ 目录。before/M6.5-B summary SHA256 分别为 3b1265b33a5f6a4ff740babf6b5b9452987da738c24799255ab2b67973fae824 和 9c58c3e39719c1a95b567266f7e316c4e255287a62617dc37816ac247f6988b0。本阶段没有改变 prefill 的单序列执行；ragged/chunked prefill batching 留给 M6.5-C，M7 可将已验证的 packed decode 路径回流为推荐实现。

M6.5-C 在 commit 27bb68c33062d64b33f7c9a578099ea52dc72251 实现真实 ragged chunked prefill batching。Scheduler 的稳定顺序被组装成同一个 PackedPagedBatch；Runner 对 token-major packed rows 只执行一次模型图、一次 logits D2H 和一次最终 synchronize， HIP PagedPrefillGqaTrusted 根据每行 offsets、past length 和 block table 隔离 KV 访问。模型图成功后才原子推进整个 prefill batch 的 KV，逐行 sampling、failure isolation、 Scheduler apply 和事件发布仍维持原有顺序。prefill 与 decode 继续使用两个独立图和 failure domain，decode 仍走专用 PagedDecodeGqaTrusted。

M6.5-B baseline commit 为 bdea0ccff817245c0fc66d1ae9f30a9a2a23e67d。相同模型 revision、device、Release flags、server limits 和五轮场景下，M6.5-C candidate 的实测结果为：

场景	指标	M6.5-B	M6.5-C	变化
batch-mixed, C4	TTFT p50	1286.229 ms	1282.168 ms	改善 0.32%
batch-mixed, C4	prompt throughput	483.767 token/s	484.735 token/s	1.002x
batch-decode, C4	completion throughput	31.440 token/s	35.331 token/s	1.124x
latency-prefill, C1	TTFT p95	302.670 ms	313.379 ms	+3.54%

correctness 和 service evidence 通过：prefill_batch_calls=14、 max_prefill_batch_size=4、max_prefill_batch_tokens=64，同时保持 decode_batch_calls=15、max_decode_batch_size=4；HIP allocation calls 为 3 -> 3， KV high-water 为 16 blocks 且最终全部回收。Host 115 tests、Python tooling 96 tests 加 21 subtests、HIP backend 7 tests、HIP operators 9 tests、Tiny 4 tests 和 Engine 9 tests 均通过。

M6.5-C performance Gate 未通过。TTFT 改善要求至少 30%，prompt throughput 要求至少 1.5x，实际分别只有 0.32% 和 1.002x；两项逐轮方向一致性均为 3/5，没有达到 4/5。 decode throughput 为 1.124x，超过不低于 0.9x 的门槛；全部场景 p95 最大回退为 latency-prefill TTFT 的 +3.54%，低于 10% 上限。额外 C4 诊断还观察到请求没有在首个 tick 全部合齐：该 workload 的 /metrics 最大 prefill batch 为 3 rows/768 tokens，首条请求 TTFT 约 300 ms，其余三条约 1.25--1.31 s。

独立 Device kernel timeline 显示 hipBLAS GEMM family 占总 kernel 时间 56.59%， PagedPrefillScoresKernel 占 29.34%，PagedPrefillOutputKernel 占 3.73%。这说明 ragged graph 已减少逐请求 launch、D2H 和 synchronize 边界，但没有减少 token 级 GEMM 或 causal attention FLOPs；在 256-token chunks 上，这两部分已经主导 TTFT。admission coalescing、prefill attention fusion、prefix cache 或 unified mixed graph 只作为后续独立设计和测量项，本阶段不以未验证优化覆盖失败证据。

Git 忽略的 M6.5-C 证据位于 build-perf/m65c-candidate-summary.json、 build-perf/m65c-service-inspection.json、build-perf/m65c-chunked-inspection.json、 build-perf/m65c-evidence.json 和 build-perf/m65-runs/candidate-round-*.json。baseline/ candidate summary SHA256 分别为 9c58c3e39719c1a95b567266f7e316c4e255287a62617dc37816ac247f6988b0 和 602a3de3ca3776c7a6d288ca4e4efebde36dd326ef98f4397e73e3e3e2fee391；最终 evidence 保留 correctness.passed=true、comparison.passed=false 和 passed=false。

M6.5-D 在 commit d693ec7f4756e4d989217272cb748d6825af9e5c 将优化重心从调度转到 prefill 主导算子。新增的独立 operator microbenchmark 固定 gfx906、20 次 warmup 和 100 次计时迭代，并同时验证数值误差、steady-state allocation 和融合前后耗时。HIP Runner 只有在 microbenchmark 门槛通过且 28 层权重均满足连续布局时才启用对应路径：

Paged prefill 先用既有 append kernel 写入 K/V pool，再用单个 wave64 online-softmax kernel 直接生成 attention output，不再运行 materialized score/output 两个 kernel；
Q/K/V 权重在模型 materialization 时连续排列，一次 hipBLAS GEMM 生成 packed QKV，再由 split kernel 写回现有 Q、K、V workspace；
Gate/Up 权重同样连续排列，一次 hipBLAS GEMM 生成 packed Gate/Up，再由单个 kernel 完成 SwiGLU 并写入 Down projection 的输入。

CPU 和不能保持 256-byte 连续权重组的 Tiny fixture 继续使用原路径；base-class 新 virtual method 保留显式 Unsupported，不会静默伪装成融合实现。Decode attention 仍使用 PagedDecodeGqaTrusted，但通过门槛的 QKV 和 Gate/Up fusion 同时用于 prefill 与 decode。为保持旧的非 batch/兼容 attention 入口，当前 Runner 仍保留完整 FP32 attention_scores slot；默认服务 max_batched_tokens=2048 下，移除 12 MiB up、增加 16 MiB packed QKV 和 24 MiB packed Gate/Up，workspace 净增加约 28 MiB。后端 allocation calls 保持 3 -> 3。

正式 operator artifact 的代表结果如下，全部 case correctness 与 allocation stability 均通过：

算子/场景	Baseline	Fused	加速比
Attention single-64	0.313 ms	0.131 ms	2.385x
Attention single-256	7.454 ms	1.148 ms	6.490x
Attention single-1024	98.565 ms	15.148 ms	6.507x
Attention equal-4x64	0.977 ms	0.350 ms	2.788x
Attention equal-3x256	23.484 ms	3.029 ms	7.752x
Attention ragged-4	16.216 ms	3.822 ms	4.243x
QKV，token-weighted geometric mean	-	-	1.132x
Gate/Up，token-weighted geometric mean	-	-	1.059x

三类融合路径因此全部被 Runner 选中。operator JSON 位于 build-perf/m65d-operator-benchmark.json，SHA256 为 de5402eb6af9ee1d9f10fdc8c6d647bbcd6ecab1f6afcb8c31d6bb61cd224731。

端到端比较使用不可变 M6.5-C candidate commit 27bb68c33062d64b33f7c9a578099ea52dc72251 作为 baseline，在相同模型 revision、 device、Release flags、server limits 和六场景矩阵下重新运行五轮：

场景	指标	M6.5-C	M6.5-D	变化
batch-mixed, C4	TTFT p50	1282.168 ms	475.622 ms	改善 62.90%
batch-mixed, C4	prompt throughput	484.735 token/s	884.815 token/s	1.825x
batch-decode, C4	completion throughput	35.331 token/s	44.916 token/s	1.271x

batch-mixed TTFT 和 prompt throughput 的逐轮方向均为 5/5。除 output-token 为 1、两侧 TPOT 都为 0 的 long-prefill 外，所有代表场景 TTFT/TPOT p95 都得到改善；最小改善是 batch-mixed TPOT p95 的 17.50%，因此 p95_no_regression=true。Service evidence 记录 prefill_batch_calls=14、max_prefill_batch_size=4、 max_prefill_batch_tokens=64、decode_batch_calls=15、 max_decode_batch_size=4，KV high-water 为 16 blocks 且最终全部回收。

baseline/candidate summary SHA256 分别为 602a3de3ca3776c7a6d288ca4e4efebde36dd326ef98f4397e73e3e3e2fee391 和 69f0ae288489e1f4522d3c0f16658627b2c68cf7d940816b40a7a9f6f11800d7。最终 build-perf/m65d-evidence.json SHA256 为 b6e3f855c7870f62a07bcb8c532c5bdda419b8119ffb63f5aa9712deae9a832f，记录 certification=true、correctness、operator selection、五轮 comparison 与总 passed=true。正式五轮按协议关闭 profiler；M6.5-C 的独立 timeline 只用于确定优化方向，本阶段不把它冒充为 M6.5-D 的 kernel attribution 证据。

M6.5-E Decode Tail 与内存优化

M6.5-E 的正式服务性能 candidate 为 commit 4fb9253eac18bf89a4203d8d5d8bcbdaff799a33，不可变 baseline 仍为 M6.5-D commit d693ec7f4756e4d989217272cb748d6825af9e5c。本阶段保留全部旧算子作为 oracle/fallback，用同一份 gfx906 MeasuredOperatorPolicy 独立选择 Fused Decode Attention、DeviceGreedy、 Q/K Norm+RoPE 和 Add+RMSNorm。正式 operator benchmark 固定 20 次 warmup、100 次计时；所有 case 的数值、row/token parity 和 steady-state allocation 检查均通过。

M6.5-E 子门	加权几何均值	生产策略	结果
Fused Paged Decode Attention	5.277457x	`true`	12/12 correctness/allocation 通过
DeviceGreedy LM-head tail	39.979739x	`true`	3/3 通过；B1/B4/B16 为 6.771x/19.302x/53.592x
Fused Q/K RMSNorm+RoPE	2.334533x	`true`	T1--T768 共 6/6 通过
Add+RMSNorm	1.402338x	`true`	T1--T768 共 6/6 通过
Canonical packed projection	1.048404x	`false`	未达到 1.05x 门槛
In-place Gate/Up	0.993444x	`false`	未达到 1.05x 门槛

后两项只关闭新增 canonical/in-place 路由；M6.5-D 已选中的 packed QKV 与 packed Gate/Up GEMM 仍保留。也就是说，当前生产路径仍运行 SplitQkvKernel 和独立 packed Gate/Up activation，不能把这两项描述为已消除。

服务 workspace 切换为 BatchService 后，legacy 的 262,144-byte FP32 attention score slot 被删除：单 Runner workspace 从 5,464,320 bytes 降为 5,202,176 bytes， attention_score_bytes=0。真实 service inspection 为 allocation calls 4 -> 4、 decode_batch_calls=15、max_decode_batch_size=4、prefill_batch_calls=14、 max_prefill_batch_size=4、KV high-water 16 blocks 且最终全部回收。Chunked inspection 另行验证 128/256-token chunk 等价、decode/prefill 并行推进、allocation calls 3 -> 3 和 KV 全回收。

最终 timeline driver 在 commit 9c5fee9 明确使用生产选择的 RunnerWorkspaceMode::kBatchService 与 PagedBatchOutputMode::kDeviceGreedy。这是必要的证据修正：早期 after trace 仍走 legacy/HostLogits，不能证明 E2/E3/E4。修正后的 15 个固定 rocTX range 全部出现 FusedQkRmsNormRopeKernel、AddRmsNormKernel 和两阶段 BatchedArgmax，12 个 decode range 出现 FusedPagedDecodeAttentionKernel；旧 PagedDecodeScoresKernel 与 PagedDecodeOutputKernel 均为 0 次。

Timeline 汇总（15 ranges）	M6.5-D E0	M6.5-E selected	变化
Device time	3,270,195.450 us	2,013,961.856 us	-38.41%
Kernel calls	7,119	5,133	-27.90%
Attention time/calls	1,977,560.093 us / 1,176	736,670.088 us / 840	-62.75% / -336 calls
GEMM time/calls	1,237,583.699 us / 1,710	1,238,349.490 us / 1,710	+0.06% / 不变
Elementwise time/calls	55,051.658 us / 4,233	38,509.798 us / 2,553	-30.05% / -1,680 calls
Device tail time/calls	0 us / 0	432.480 us / 30	新增 Device argmax

旧 decode score/output kernel 合计 672 次、1,515,105.211 us，被 336 次、 283,393.109 us 的 fused decode attention 替代。E4 在全部 range 中分别记录 420 次 Fused Q/K kernel（12,024.651 us）和 420 次 Add+RMSNorm kernel（7,227.687 us）。单个 decode range 的 kernel calls 从 480 降至 342；单个 prefill range 从 453 降至 343。代表性的 B4 decode Device time 改善随 context 从 L64 的 6.31%、L256 的 18.74%、 L1024 的 42.51% 增长到 L4096 的 61.46%；B16/L4096 改善 76.98%。Prefill L64/L256/ L1024 分别改善 2.65%/2.54%/1.93%，符合本阶段优化主要作用于 decode 的预期。

正式端到端比较关闭 profiler，按相同模型 revision、device 0、Release flags、server limits 和六场景协议重新运行五轮。下表均为五轮中位数：

场景	TTFT p50，D -> E	TPOT p50，D -> E	Completion throughput，D -> E
batch-decode, C4	217.392 -> 208.075 ms	80.383 -> 74.640 ms	44.916 -> 48.164 token/s
batch-mixed, C4	475.622 -> 454.806 ms	97.560 -> 77.984 ms	27.650 -> 31.930 token/s
latency-prefill, C1	113.712 -> 110.371 ms	89.815 -> 77.712 ms	10.770 -> 12.221 token/s
latency-short, C1	80.296 -> 76.904 ms	78.540 -> 74.333 ms	12.690 -> 13.389 token/s
long-prefill, C1	785.672 -> 764.197 ms	0 -> 0 ms	1.272 -> 1.307 token/s
saturation, C16	297.271 -> 279.964 ms	86.919 -> 75.237 ms	140.346 -> 157.349 token/s

场景	TTFT p95，D -> E	TPOT p95，D -> E	p95 结论
batch-decode	218.399 -> 208.963 ms	80.885 -> 74.867 ms	均改善
batch-mixed	476.733 -> 455.688 ms	100.838 -> 78.247 ms	均改善
latency-prefill	114.266 -> 110.526 ms	90.457 -> 77.792 ms	均改善
latency-short	80.473 -> 77.608 ms	78.746 -> 74.404 ms	均改善
long-prefill	786.962 -> 765.593 ms	0 -> 0 ms	TTFT 改善，TPOT 不适用
saturation	302.131 -> 285.241 ms	88.250 -> 76.468 ms	均改善

四个 primary direction 都是 5/5：C4 decode TPOT、C4 decode throughput、C16 throughput 和 C4 mixed TTFT 每轮方向均改善；p95_no_regression=true。但总性能 certification FAILED，且不放宽阈值：C4 decode TPOT 只改善 7.1448%，未达到 >=10%；C4 decode throughput 只有 1.072299x，未达到 >=1.10x。C16 throughput 为 1.121148x、C4 mixed TTFT ratio 为 0.956232x，这两项通过。最终 evidence 因此保留 correctness.passed=true、comparison.passed=false 和总 passed=false。

完整门禁记录 Host 117/117、HIP ctest 117/117、tooling 116 passed + 50 subtests、 HIP backend 8/8、HIP operators 19/19、Tiny 6/6 和 Engine 10/10。M4B、M5 与 M6 quick 回归也分别通过；verify_m65e.sh 的退出码 1 只表示上述硬性能认证失败，不是正确性或测试基础设施失败。

Git 忽略的正式 artifact SHA256 为：M6.5-D baseline summary 69f0ae288489e1f4522d3c0f16658627b2c68cf7d940816b40a7a9f6f11800d7，E0 timeline d740faa240d33138281daf5ee941b4934810f244295932cc4cc567df23eac1d0，M6.5-E operator benchmark 040fce2792ea8589ba1a34e5d6fdd7752924569299f10eca9daf2de4524f4dd5，after timeline 047838f6fa2000db0c7ce68535293edbb2ab1315abb4427aaeb052dbfee924d3，五轮 summary 74c1f072bf18725d04e41eb6c6e2838bfc6b25e3bfaffd9e81022b6d1463e15a，最终 evidence 277377df160173f20001d7d054a29a16ff7ff85eb0b2b2f03c5213ddccbe8d27。

M6.5-F Device Sampling 与 LM Head 候选

M6.5-F 把 stochastic sampling 从 HostLogits fallback 移到 HIP：Runner 创建时一次性预分配 double sort keys、token IDs、segment offsets 和 hipCUB scratch；每步只上传 per-row sampling config、counter 与排序去重后的 history，最后只 D2H token ID/validity。 simple greedy 继续走原两阶段 BatchedArgmax。Device 路径保持 Host SampleToken 的 repetition penalty、temperature、top-k、top-p、SplitMix64 counter RNG、stable tie 和 NaN/Inf 语义，混合 greedy/stochastic batch 不再回传完整 logits。

固定 Qwen3-0.6B shape（hidden 1024、vocab 151936）的 5 次 warmup/20 次计时结果如下。所有 case 的 token/RNG parity 与 steady-state allocation 检查均通过：

Case	Host Sampling -> Device Sampling	Speedup
B1	22.422 -> 12.337 ms	1.818x
B4	89.129 -> 12.544 ms	7.106x
B16	353.525 -> 12.613 ms	28.030x

Device Sampling 加权几何均值为 18.945717x，且没有 case 低于 0.95x，因此 MeasuredOperatorPolicy::kDeviceSampling=true。kernel-only rocprof timeline 共 142 个 dispatch；sampling 的 segmented radix sort 为 6 calls/75,315.988 us，preprocess 为 6/238.716 us，select 为 6/295.680 us，row init 为 6/46.400 us。也就是说，目前约 12.6 ms 的 Device Sampling 时间几乎全部来自全词表 radix sort，下一步应优化 top-k candidate selection，而不是继续压缩 config/preprocess kernel。预分配 sampling workspace 在 B1/B4/B16 分别为 5,469,952/21,879,040/87,515,392 bytes；真实 B4 service inspection 的总 Runner workspace 为 27,084,032 bytes，且 allocation calls 4 -> 4、KV 全回收。

同时实现了不落 full logits 的 FusedLmHeadArgmax 候选，并与 hipBLAS Linear + BatchedArgmax 比较：

Case	hipBLAS+Argmax -> Fused LM Head	Speedup
B1	2.809 -> 3.920 ms	0.717x
B4	2.827 -> 6.267 ms	0.451x
B16	2.910 -> 17.242 ms	0.169x

候选加权几何均值仅 0.218025x，三个 case 全部低于 0.95x，故保留实现和 oracle 测试，但生产策略为 kFusedLmHeadArgmax=false。这说明该卡上省掉 logits 写回的收益远小于自定义 LM Head 丢失 hipBLAS GEMM 效率的代价。

默认 greedy 服务另按 M6.5-E 相同协议关闭 profiler 重跑五轮。相对 M6.5-E selected summary，六场景最差变化为 latency-short TTFT p50 76.904 -> 77.709 ms（+1.05%）；其余 TTFT/TPOT p50/p95 和 completion throughput 偏差均不超过约 0.25%，没有性能回退。 M6.5-F operator benchmark、greedy 五轮 summary 和 kernel CSV 的 SHA256 分别为 d323869610e64fc38fbe1721abe363fec0f029ca735f463c1f8781c557d84d34、 58bec33bc0d2ed85b7785c15ab5ac7503e098271e12a094b7fbaeaf47776ea1c 和 8517d2fd14009fcf0202ced6de61421b992d413e621f88f12e2d32ebbbdf75ca。

Qwen3-4B 单卡 Bring-up

在 M6.5-F 单卡基线之上，运行时增加了第二个严格白名单 profile：Qwen3-4B。官方 revision 为 1cfa9a7208912126459214e8b04321603b3df60c，对应配置为 hidden_size=2560、intermediate_size=9728、36 层、32 个 Query Head、8 个 KV Head、head_dim=128。权重清单固定在 model_manifest_qwen3_4b.json，3 个 Safetensors shard 的总 BF16 字节数为 8,044,936,192。

4B checkpoint 的 index 没有独立 lm_head.weight，因为配置启用了 tied word embeddings。WeightArchive 现在允许这个唯一的可选张量，校验后把 LM Head 设备视图直接别名到 model.embed_tokens.weight，不会在显存中复制 embedding。

同步和离线校验：

python3 tools/sync_model.py sync \
  --manifest model_manifest_qwen3_4b.json \
  --dest "$HOME/.cache/metainfer/models/Qwen3-4B" \
  --endpoint https://hf-mirror.com
python3 tools/sync_model.py check \
  --manifest model_manifest_qwen3_4b.json \
  --dest "$HOME/.cache/metainfer/models/Qwen3-4B" \
  --endpoint https://unreachable.invalid

在当前 gfx906/Z100SM 单卡上使用保守容量启动：

build-perf/bin/metainfer_ref_qwen3_server \
  --model-dir "$HOME/.cache/metainfer/models/Qwen3-4B" \
  --hip-device 0 --port 18004 \
  --max-context-tokens 1024 --max-output-tokens 32 \
  --max-active-sequences 1 --max-queue-size 8 \
  --max-batched-tokens 64 --prefill-chunk-tokens 64

实机结果（2026-07-21，profiler disabled）：

项目	结果
服务启动后稳态显存	8,472--8,474 MiB
GPU 数/单卡显存	4 / 16,368 MiB
真实请求完成	5/5，failed 0，cancelled 0
最大 Prefill batch	64 tokens
101-token prompt	成功拆为 64+37 两个 chunk
8-token greedy TTFT	440.7 ms
greedy 相邻 token 间隔中位数	395.7 ms
stochastic temperature=0.7	成功，Device Sampling 路径
请求结束 KV blocks	0
优雅退出后 GPU 显存	2 MiB

同一 chat prompt 的 CPU Transformers 首 token 为 ID 151667（<think>），原生 C++ 服务流式首 token 相同。该实验证明 4B 已经可以在单卡完整运行，但尚未替代 0.6B 的完整 M3 golden/e2e certification；下节补充五轮性能基线，numerical golden 仍需在进入 TP2/TP4 前单独生成。

Qwen3-4B 五轮正式性能基线

4B baseline 使用独立 Qwen3-4B/baseline schema role。Profiler 关闭，Release HIP flags 包含 -O3 -DNDEBUG --offload-arch=gfx906，同一服务进程连续运行 5 轮；每个场景每轮先执行 2 个不计入结果的 warmup request，再测量 4 个 request。单卡显存限制下将 0.6B 的 C16 saturation 替换为 C4，其他场景 token 规格保持一致。

场景	TTFT p50/p95	TPOT p50/p95	Prompt tok/s	Completion tok/s
latency-short C1	398.967 / 399.118 ms	396.051 / 396.192 ms	10.087	2.522
latency-prefill C1	581.581 / 583.432 ms	400.389 / 400.556 ms	75.656	2.364
batch-decode C4	1050.302 / 1051.501 ms	396.531 / 396.736 ms	18.287	9.143
batch-mixed C4	2692.954 / 2693.455 ms	401.009 / 401.286 ms	186.141	5.817
saturation C4	1050.247 / 1050.915 ms	396.405 / 396.777 ms	33.452	8.363
long-prefill C1	3110.714 / 3112.192 ms	不适用	329.124	0.321

五轮中，主要 TPOT 和 completion throughput 的 (max-min)/median 均不超过 0.18%。相对 M6.5-F 的 0.6B 同场景，4B Decode TPOT 为 5.14--5.33x，C1/C4 completion throughput 为 0.182--0.194x；4B/0.6B checkpoint 字节比为 5.35x，说明 Decode 基本随模型规模线性变化。Long Prefill TTFT 为 4.07x，大矩阵在该卡上的利用率优于 0.6B 小矩阵。

正式 summary 位于 build-perf/qwen3-4b/five-round-summary.json，五个原始轮次位于 build-perf/qwen3-4b/qwen3-4b-runs/。Summary SHA256 为 c8acf65d370259f2c642953e701e3959fc30775157907d638de32101f6052329。运行结束后服务 SIGTERM 退出耗时 2.617 s，GPU 0 显存回到 2 MiB。

本次遵循正式五轮测量协议，但运行于包含 M6.5-F/4B 未提交修改的 feature worktree，因此属于 working-tree baseline，不是 immutable release certification。冻结源码提交后需要再运行一次，才能把 commit 字段作为可复现的唯一源码身份。

Qwen3-4B TP2 Bring-up

4B TP2 第一阶段使用单进程双 Rank，并固定在 HIP device 0/1：

Rank 0: Backend 0 + Stream 0 + Operators 0 + local weights + local KV
Rank 1: Backend 1 + Stream 1 + Operators 1 + local weights + local KV

WeightArchive::MaterializeTensorParallel 按以下规则物化每个 Rank：Q/K/V 和 Gate/Up 按输出行切分；O 和 Down 按输入列切分，并在 Host 端打包连续列后执行一次 H2D；Embedding、RMSNorm、Final Norm 和 tied LM Head 复制。每个 Rank 的局部 shape 为 16 个 Query Head、4 个 KV Head、intermediate_size=4864，KV Cache 也只保存本地 4 个 KV Head。每卡权重 allocation 为 4,411,620,352 bytes，约为单卡完整 checkpoint 的 54.8%，其中复制的 Embedding/LM Head 限制了继续下降的空间。

每层在 O Projection 后、Down Projection 后各执行一次 BF16 AllReduce。当前无 RCCL， Collective 使用双向 HIP P2P：两个 Stream 先把 partial hidden 写入本地只读 scratch，通过跨设备 HIP event 建立依赖，再由两个 reduction kernel 同时读取本地和对端 scratch，结果写回各自 hidden。该设计避免 CPU staging，也避免双方直接原地求和产生读写竞争。 metainfer_ref_device_probe --require-peer 0:1 已实测四张卡之间的全部有向 P2P 路径均可启用。

构建和 correctness inspection：

cmake --build build-perf --target \
  metainfer_ref_tp2_inspect metainfer_ref_tp2_collective_benchmark -j4
build-perf/bin/metainfer_ref_device_probe \
  --min-devices 2 --require-peer 0:1
build-perf/bin/metainfer_ref_tp2_inspect \
  --model-dir "$HOME/.cache/metainfer/models/Qwen3-4B" \
  --devices 0,1 --max-new-tokens 8 --batch-size 1

同一 32-token chat prompt 下，TP2 两个 Rank 连续五轮都生成：

[151667, 198, 99692, 3837, 20002, 56007, 100146, 2073]
<think>\n好的，用户问的是“

该文本与 4B 单卡 HTTP 服务的 8-token greedy 输出一致；每一步每 Rank 的 AllReduce 次数都严格等于 36 layers * 2，没有 Rank token 分歧。

TP2 五轮性能结果

Release/gfx906、profiler disabled 的 direct-runner 五轮结果如下：

场景	TTFT p50/p95	TPOT p50/p95	Completion tok/s	对单卡参考
C1, 32 prompt, 8 output	245.026 / 250.067 ms	224.026 / 224.163 ms	4.464	TTFT 1.63x, TPOT 1.77x
C4, 32 prompt, 8 output	201.866 / 203.269 ms	224.611 / 224.824 ms	17.809	TPOT 1.77x, throughput 1.95x

C1/C4 分别包含 35 个 decode timing sample，最大和最小 TPOT 相差不到 0.34 ms。单卡参考来自 qwen3-4b/five-round-summary.json 的相同 token shape，但经过 HTTP 和 Scheduler；TP2 数据直接调用 Runner。因此 C1 的 TTFT 比值只作为最近似 shape 参考， C4 的 1050 ms -> 202 ms 不计算加速比，不能归因给 TP2 本身。TPOT 和 steady decode throughput 受 HTTP 的影响较小，是当前更可信的比较。

P2P Collective 使用 50 次 warmup、1000 次计时：

Hidden payload	AllReduce latency	Payload bandwidth
Decode B1, 5,120 bytes	57.301 us	0.089 GB/s
Decode B4, 20,480 bytes	58.966 us	0.347 GB/s
Prefill 32, 163,840 bytes	75.041 us	2.183 GB/s
Prefill 256, 1,310,720 bytes	342.891 us	3.823 GB/s

B1 Decode 每 token 的 72 次 Collective 估算约 4.13 ms，仅占 224.03 ms TPOT 的约 1.8%。当前剩余时间主要来自本地半尺寸 GEMM，以及两个 Rank 重复执行的 Norm、 Attention、LM Head 和 Argmax，而不是 PCIe P2P。

正式 artifacts：

build-perf/qwen3-4b-tp2/five-round-summary.json，SHA256 f93c60d0f1466701b01f97ef973e3674d209ce0079a7c4c295508751106db914；
build-perf/qwen3-4b-tp2/c4-five-round-summary.json，SHA256 b6224682e25f5636520d3d12c9198685c189a0b6b245f596540d2828a8b61145；
build-perf/qwen3-4b-tp2/collective-benchmark.json，SHA256 48396d9722fa30e5d2f96f18793a9b265f1958fa6825687391f7d90d49eaa689。

TP2 HTTP/OpenAI 服务

TP2 已接入既有 ServiceRuntime 和 OpenAiHttpServer。同一个 server 可通过 --hip-device N 启动单卡，或通过 --devices FIRST,SECOND 启动单进程双 Rank：

build-perf/bin/metainfer_ref_qwen3_server \
  --model-dir "$HOME/.cache/metainfer/models/Qwen3-4B" \
  --devices 0,1 --host 127.0.0.1 --port 18008 \
  --max-context-tokens 512 --max-output-tokens 32 \
  --max-active-sequences 4 --max-queue-size 16 \
  --max-batched-tokens 256 --prefill-chunk-tokens 128 \
  --block-size-tokens 16

启动日志会打印 tensor_parallel_world_size=2 devices=0,1。HTTP 层不需要 TP2 专用协议，以下端点保持单卡服务兼容：

GET /health
GET /v1/models
POST /v1/completions
POST /v1/chat/completions，包括 SSE stream=true
GET /metrics

真实 Qwen3-4B TP2 服务验证已覆盖模型发现、非流式 completion、流式 chat completion、双并发请求、客户端中途断连和 SIGTERM。断连请求会由 HTTP chunk provider 触发双 Rank Cancel，服务继续可用；SIGTERM 后 GPU 0/1 显存均回到约 2 MiB。增量 tokenizer 还会缓存不完整 UTF-8 尾部，生成在半个字符处结束时输出 U+FFFD，避免 JSON/SSE 进程崩溃。

该 TP2 阶段当时只支持 world_size=2 和 P2P 可达的 HIP device；每个请求由两个 Rank 接收并并发执行，只有 Rank 0 执行完整 LM Head/Sampling，再向 Rank 1 广播 token 和 valid。Collective 已支持 group-wide abort，任一 Rank 在进入同步点前失败时会唤醒对端并让服务统一失败退出。当前版本已经进一步加入 TP4、RCCL 和 vocab-parallel LM Head，实现与验证见文末的 TP4 章节；多进程 Rank 启动仍未加入。

Qwen3-8B TP2 Bring-up

8B 使用固定官方 revision b968826d9c46dd6066d109eabc6255188de91218。安全清单包含 11 个文件、逐文件大小和 SHA256：

python3 tools/sync_model.py sync \
  --manifest model_manifest_qwen3_8b.json \
  --dest "$HOME/.cache/metainfer/models/Qwen3-8B" \
  --endpoint https://hf-mirror.com
python3 tools/sync_model.py check \
  --manifest model_manifest_qwen3_8b.json \
  --dest "$HOME/.cache/metainfer/models/Qwen3-8B"

8B profile 为 hidden_size=4096、intermediate_size=12288、36 layers、32 Query Heads、8 KV Heads 和 head_dim=128。TP2 Rank-local shape 为 16 Query Heads、4 KV Heads 和 intermediate_size=6144。与 4B 不同，官方 8B 使用 tie_word_embeddings=false，Embedding 和 LM Head 是两个独立的 [151936,4096] 权重。初始基线在每个 Rank 都复制两者，每 Rank 实际权重 allocation 为 9,435,703,296 bytes，约 8.79 GiB；当前实现只在 Rank 0 物化 LM Head，Rank 1 权重降为 8,191,043,584 bytes。

真实模型 correctness inspection：

build-perf/bin/metainfer_ref_tp2_inspect \
  --model-dir "$HOME/.cache/metainfer/models/Qwen3-8B" \
  --devices 0,1 --max-new-tokens 8 --batch-size 1 \
  --expected-first-token 151667

五轮均生成：

[151667, 198, 99692, 3837, 20002, 56007, 100146, 2073]
<think>\n好的，用户问的是“

每一步每 Rank 都执行 36 * 2 = 72 次 AllReduce，Rank token、Collective 次数和 first-token golden 全部一致。额外使用单卡 device 2、64-token context、单 active sequence 跑过同一 32-token Prompt，前两个输出同为 [151667,198]；该单卡路径仅作为分片正确性对照，不作为 8B 单卡性能基线。

Release/gfx906、profiler disabled、fresh process per round 的优化前五轮结果：

场景	TTFT p50/p95	TPOT p50/p95	Completion tok/s	每 Rank allocation high-water
B1, 32 prompt, 8 output	357.577 / 358.702 ms	335.792 / 336.069 ms	2.978	9,447,825,664 bytes
B4, 32 prompt, 8 output	329.228 / 331.138 ms	336.507 / 336.882 ms	11.887	9,507,771,392 bytes

相对同方法的 4B TP2：8B 每 Rank 权重为 2.139x，B1/C4 TPOT 都约为 1.499x， token throughput 约为 4B 的 66.7%。另一方面，8B B4 completion throughput 是 B1 的 3.992x，而 TPOT 只增加约 0.21%，说明 Packed Decode/Continuous Batching 在该 shape 上仍保持接近线性的 batch 扩展。8B 与 4B 的 layers、Q/KV Head 数相同，因此 KV bytes 在相同服务容量下相同；主要增量来自更大的 Hidden/MLP GEMM、1.6 倍的 AllReduce payload 以及更大的独立 LM Head。

正式 artifacts：

build-perf/qwen3-8b-tp2/five-round-summary.json，SHA256 5e7cabeba22cde14474df57da6df77e4f9322a53f2cc331e878bb313b79c38bb；
build-perf/qwen3-8b-tp2/c4-five-round-summary.json，SHA256 9efc9e9cb125b516b41098d74c6f5c7882e2d0cdb158c5022fd8867266831e61。

Rank 0-only LM Head/Sampling

TP2 BatchService 当前只在 Rank 0 执行 Final RMSNorm、完整 LM Head 和 Device Greedy/Sampling。Rank 1 不物化 untied lm_head.weight，也不分配 logits、argmax partial 和 sampling workspace；两 Rank 在每一步 Transformer AllReduce 完成后调用两次 P2P broadcast，同步 batch_tokens 和 batch_valid，因此仍向 Scheduler 返回完全一致的 Rank-local output。Collective descriptor 同时校验 operation、root、dtype 和 bytes，避免 AllReduce/Broadcast 调用顺序不一致时被误配。

同一 B1、2-token trace 的前后对比：

Kernel/设备统计	优化前	Rank 0-only	变化
LM Head GEMM	4 calls / 43.649 ms	2 calls / 21.816 ms	Rank 1 两次调用消除
Final Norm	148 calls / 4.135 ms	146 calls / 4.029 ms	Rank 1 tail norm 消除
Argmax partial/final	4 / 4 calls	2 / 2 calls	Rank 1 sampling 消除
Rank 0 kernel sum	692.985 ms	693.191 ms	基本不变
Rank 1 kernel sum	693.365 ms	668.605 ms	-24.760 ms
P2P copy kernels	288 / 0.881 ms	292 / 0.898 ms	+4 / 0.017 ms

Kernel 时间是两步推理的设备时间之和。LM Head 在两个设备上原本并行执行，因此消除 Rank 1 的副本不会缩短由 Rank 0 决定的临界路径；新增的 4 次极小 broadcast copy 可以忽略，但两次 Collective 仍有主机同步开销。五轮正式复测印证了这一判断：

场景	TTFT p50/p95	TPOT p50/p95	Completion tok/s	Rank 0/1 allocation high-water
B1, 32 prompt, 8 output	359.346 / 366.621 ms	335.799 / 336.031 ms	2.978	9,447,825,664 / 8,197,083,904 bytes
B4, 32 prompt, 8 output	332.269 / 333.361 ms	336.588 / 336.902 ms	11.884	9,507,771,392 / 8,238,790,144 bytes

相对优化前，B1 TPOT p50 为 +0.002%、吞吐为 -0.002%；B4 TPOT p50 为 +0.024%、吞吐为 -0.024%，均为持平。Rank 1 的 B1/B4 allocation high-water 分别减少 1,250,741,760 和 1,268,981,248 bytes。该变化的主要收益是容量和能耗，而不是延迟；若要让 LM Head 缩短临界路径，需要把 vocab 行切到两个 Rank 并做分布式 top-k，即 vocab parallel，而不是把完整 LM Head 移到单个 Rank。

Rank 0-only artifacts：

build-perf/qwen3-8b-tp2/root-only-five-round-summary.json，SHA256 b5d3ddbf8f1ef31ceecfca14dc5e96009171b617afd544ec8035a9583ac89de6；
build-perf/qwen3-8b-tp2/root-only-c4-five-round-summary.json，SHA256 2b415bee426fe9ae0423b6e3a5d5a785085abf01079bdbe8be3d6366065dbf4d；
build-perf/qwen3-8b-tp2/timeline-after-root-only/results_after-root-only.csv， SHA256 af40d4c8be8b6cf24bc633d59a964579a9404c71076ebebfd982925a638c6d86。

Decode small-M Linear

Kernel Timeline 显示 Decode 阶段的 Transformer Linear 长期由 small-M hipBLAS GEMM 主导。gfx906 上这些 M=1..4、K/N 很大的问题主要读取一次 BF16 权重，通用 GEMM 内核的固定开销和数据布局并不适合该 shape。因此 HIP Backend 增加了专用 small-M Linear：一个 64-thread wave 计算一个输出通道，线程沿 K 维累加并做 wave reduction； M=2..4 在同一次权重遍历中累加多行输入，从而复用权重流量。累加使用 FP32，结果写回 BF16，并支持 Down Projection 使用的非连续 input row stride。

默认路由只在 M<=4 && output=BF16 时选择专用 Kernel。Prefill 的 M>4 和输出为 F32 的 LM Head 继续使用 hipBLAS，避免影响 TTFT 与采样数值路径。独立 microbenchmark 使用 Qwen3-8B TP2 的四组 Rank-local shape，5 次 warmup、20 次计时；8 个 case 均通过 hipBLAS 对照和 allocation-stable 检查：

Projection shape	B1 hipBLAS / small-M	B1 加速	B4 hipBLAS / small-M	B4 加速
QKV `[4096 -> 3072]`	1.634 / 0.236 ms	6.92x	1.634 / 0.193 ms	8.47x
O `[2048 -> 4096]`	0.827 / 0.125 ms	6.60x	0.825 / 0.134 ms	6.14x
Gate/Up `[4096 -> 12288]`	3.921 / 0.430 ms	9.11x	3.920 / 0.415 ms	9.45x
Down `[6144 -> 4096]`，stride 12288	2.440 / 0.190 ms	12.86x	2.441 / 0.233 ms	10.47x

专用 Kernel 的有效权重带宽为 106.5-265.4 GB/s，而这些 shape 的 hipBLAS 为 15.4-25.7 GB/s。真实 8B TP2、32-token prompt、8-token output 五轮复测结果如下；基线为上一节 Rank 0-only 版本：

场景	基线 TPOT p50	small-M TPOT p50/p95	TPOT 加速	基线 / small-M throughput	small-M TTFT p50/p95
B1	335.799 ms	56.450 / 58.375 ms	5.95x	2.978 / 17.715 tok/s	359.266 / 360.513 ms
B4	336.588 ms	55.286 / 56.961 ms	6.09x	11.884 / 72.350 tok/s	331.786 / 334.062 ms

五轮输出都与优化前逐 token 一致： [151667,198,99692,3837,20002,56007,100146,2073]。每轮每 Rank 有 1008 次 small-M 调用，即 7 个 Decode step 的 4 projections * 36 layers；Rank 0/1 分别保留 152/144 次 hipBLAS 调用，来自 Prefill 和 Rank 0 LM Head。B1 两步 Kernel Timeline 中，优化前 Prefill 和 Decode 共 576 次 Transformer hipBLAS GEMM 为 1271.106 ms；优化后 288 次 Prefill hipBLAS GEMM 为 636.758 ms，288 次 Decode small-M Kernel 为 70.503 ms。两卡 kernel sum 从 1361.796 ms 降到 798.438 ms，其中 Rank 0 为 693.191 -> 410.115 ms，Rank 1 为 668.605 -> 388.323 ms。TTFT 保持在原波动范围，符合只优化 Decode 的预期。

Decode small-M artifacts：

build-perf/qwen3-8b-tp2/decode-gemm-small-m-candidate.json，SHA256 3ca6a2e86ebb675fedf07cb24a368bc07e06da060d1df2f3221004f48d11352a；
build-perf/qwen3-8b-tp2/small-m-five-round-summary.json，SHA256 f5d6dec67f21790608c349399ad0d4b16d3fa499e809df8607ee1958920b518a；
build-perf/qwen3-8b-tp2/small-m-c4-five-round-summary.json，SHA256 1328934a5d720c6e45ebd63228e47b3fa07c811c68761dbb588f570255e7fa79；
build-perf/qwen3-8b-tp2/timeline-small-m/results_small-m.csv，SHA256 7de1bc21b67a4bfe2ba35cdcdb5945ee6b65a2e38f45c553718c2e0642ec2c9e。

Decode packed small-M Linear v2

第一版 small-M Kernel 每个 lane 每轮读取一个 BF16 input/weight，并只维护一个输出通道。 v2 保持“一整个 64-thread wave 计算一个输出通道”的并行度不变，但将 K 维改为 BF16x2 packed load：每次 32-bit 读取两个相邻 BF16，展开为两个 FP32 FMA，从而将循环、地址计算和 load 指令数减半。M=1..4 仍在一次权重遍历中累加所有 batch 行，结果继续以 BF16 写回。默认路由还要求 K、input row stride 和 input/weight pointer 都满足 pair alignment；非标准 TensorView 自动退回 scalar small-M，Prefill 和 F32 LM Head 仍走 hipBLAS。

开发过程中也测量了一个 wave 同时计算 2/4 个输出通道的候选。它虽然增加 ILP，却减少了可调度 wave 数：只对 QKV B1 有明显收益，对 O/Down 和 4-output 组合均为负提升，因此已从生产实现删除。最终只保留 8 个 shape 全部正提升的 packed K 维方案。相同 5 次 warmup、20 次计时的正式 microbenchmark 如下：

Projection shape	B1 scalar / packed	B1 加速	B4 scalar / packed	B4 加速
QKV `[4096 -> 3072]`	0.241 / 0.125 ms	1.93x	0.193 / 0.143 ms	1.35x
O `[2048 -> 4096]`	0.119 / 0.094 ms	1.27x	0.135 / 0.119 ms	1.13x
Gate/Up `[4096 -> 12288]`	0.429 / 0.297 ms	1.45x	0.412 / 0.336 ms	1.22x
Down `[6144 -> 4096]`，stride 12288	0.196 / 0.141 ms	1.39x	0.234 / 0.169 ms	1.39x

8 个 case 均与 hipBLAS 对照一致且 steady-state allocation 不变。packed 有效权重带宽为 140.8-356.2 GB/s，第一版 scalar 为 104.5-256.2 GB/s。真实 8B TP2 五轮结果以上一节 scalar small-M 为基线：

场景	scalar TPOT p50/p95	packed TPOT p50/p95	TPOT 加速	scalar / packed throughput	packed TTFT p50/p95
B1	56.450 / 58.375 ms	42.949 / 44.860 ms	1.314x	17.715 / 23.284 tok/s	358.655 / 358.944 ms
B4	55.286 / 56.961 ms	46.748 / 48.682 ms	1.183x	72.350 / 85.565 tok/s	330.874 / 332.406 ms

B1/B4 completion throughput 分别增加 31.4%/18.3%；相对最初 Rank 0-only hipBLAS 版本，TPOT 累计加速达到 7.82x/7.20x。每轮两个 Rank 都记录 1008 次 packed small-M，8-token 输出继续逐 token 相同。最终 B1 两步 Timeline 中，288 次 Decode Linear 从 scalar 的 70.503 ms 降到 packed 的 47.015 ms，降低 33.3%；Prefill Transformer hipBLAS 为 636.758 -> 636.694 ms，LM Head 为 21.823 -> 21.816 ms，其他 Kernel 为 69.353 -> 69.141 ms，说明收益来自目标 Kernel。两卡 kernel sum 从 798.438 ms 降到 774.666 ms；Rank 0/1 分别为 410.115 -> 399.616 ms 和 388.323 -> 375.049 ms。

Decode packed v2 artifacts：

build-perf/qwen3-8b-tp2/decode-gemm-packed-v2-candidate.json，SHA256 53b9f841e1ddd28aeb9aaa32a4ab8484dda07e6c74609143e7fb036898dc8a4b；
build-perf/qwen3-8b-tp2/packed-v2-five-round-summary.json，SHA256 3c6edaf017ab5397fe60ddd3f8a76fccd805981128cbb6c6e257aee0d2da289b；
build-perf/qwen3-8b-tp2/packed-v2-c4-five-round-summary.json，SHA256 00ecfe6c7915954c0da366178ee60b45caf4a2a7695291ed6d1285dce9940be4；
build-perf/qwen3-8b-tp2/timeline-packed-v2-final/results_packed-v2-final.csv， SHA256 a7a2b72c401855b45dc285cfd9d1b2a7e12ac4a4ff7986c36779775d235d625d。

8B TP2 HTTP 服务必须显式限制容量。默认 64 * 4096 KV 配置在每个 Rank 约占 18 GiB，无法放入 16 GiB 设备。已验证的受控启动方式：

build-perf/bin/metainfer_ref_qwen3_server \
  --model-dir "$HOME/.cache/metainfer/models/Qwen3-8B" \
  --devices 0,1 --host 127.0.0.1 --port 18010 \
  --max-context-tokens 512 --max-output-tokens 32 \
  --max-active-sequences 2 --max-queue-size 8 \
  --max-batched-tokens 128 --prefill-chunk-tokens 128 \
  --block-size-tokens 16

真实服务验证覆盖 health、models、非流式 chat、SSE、两个并发请求、客户端断连取消、取消后的继续请求和 SIGINT。Rank 0-only tail 复测中，非流式和 SSE 都生成同一 <think>\n，服务运行时 Rank 0/1 分别使用约 9737/8537 MiB，退出后四张卡都回到约 2 MiB。下一性能重点应转向 vocab-parallel LM Head 或占主要时间的 Transformer GEMM，而不是继续优化 token broadcast。

BF16 TP2 性能收尾

收尾阶段固定 gfx906、BF16、Qwen3-8B/4B TP2，不引入量化、推测解码或 TP4。最终生产路由保留 vocab-parallel LM Head、实测 medium-M Prefill Linear、BF16x2 small-M 和 8B B1 的 BF16x4 small-M。BF16x4 以一次 64-bit load 读取 4 个 BF16，保持一个 wave 计算一个输出通道、FP32 累加和 BF16 输出语义。5 次 warmup、20 次计时中，B1 的 O、Gate/Up、 Down 相对 BF16x2 分别为 1.14x/1.11x/1.12x，因此只启用这三个 shape；B1 QKV 的 1.05x 余量不足，B4 除 O 外为负提升，而 B4 O 的端到端收益落入噪声，均保持 BF16x2。

两个正确但未进入生产的候选也保留了独立证据：

AllReduceSumAdd 和 AllReduceSumAddRmsNorm 在双卡测试中逐位匹配旧路径，但 fused RMS kernel 降低远端 reduce 并行度，并将 RMSNorm 纳入 rank 间 done barrier。8B B1 五轮 TPOT p50 从关闭时的 36.023 ms 退化到 36.814 ms，所以 kTensorParallelResidualFusion=false。
top_k<=64 候选保持 host token/RNG 语义，但 K=50 的重复词表扫描慢于 radix sort。 B1/B4/B16 分别为 13.115/19.266/19.447 ms，原路径为 12.345/12.554/12.617 ms，selection 明确为 false。全词表 top-p/mixed batch 继续走原 segmented radix sort。

真实模型使用同一 32-token prompt、8-token output、fresh process per round。最终五轮：

模型/场景	TTFT p50/p95	TPOT p50/p95	Completion throughput
Qwen3-8B TP2 B1	343.857 / 346.312 ms	36.023 / 36.419 ms	27.760 token/s
Qwen3-8B TP2 B4	326.851 / 328.255 ms	41.753 / 44.155 ms	95.801 token/s
Qwen3-4B TP2 B1	152.882 / 158.921 ms	19.891 / 20.359 ms	50.273 token/s

相对 Decode packed v2，8B B1 TPOT p50 降低 16.13%、throughput 增加 19.22%； B4 TPOT p50 降低 10.68%、throughput 增加 11.96%。五轮都生成完全相同的 token： [151667,198,99692,3837,20002,56007,100146,2073]。B1 每 Rank 命中 756 次 BF16x4，B4 和 4B 均为 0，证明 shape policy 没有外溢。

最终两步 Kernel Timeline 的两卡 kernel sum 为 755.756 ms，packed v2 为 774.666 ms。Decode Linear 从 47.015 ms 降到 BF16x4+BF16x2 的 43.405 ms； Prefill Linear 由 216 次 hipBLAS 577.708 ms 和 72 次 medium-M 40.940 ms 组成，合计 618.648 ms，此前为 636.694 ms。当前主要剩余热点仍是 Prefill Linear，其次是 Decode small-M Linear 43.405 ms 和 TP2 AllReduce 9.852 ms。

最终验证包括 Host CTest 134/134、HIP Backend 9/9、HIP Operators 21/21、 Tiny HIP 6/6、Dynamic Engine HIP 10/10、TP2 Collective 9/9。双卡 HTTP 同时通过 greedy completion 和 temperature=0.7/top_k=50/top_p=0.9/seed=123 chat 请求，退出后四卡显存归零。tooling 的标准库 unittest 通过 114 项；当前系统没有安装 pytest，因此唯一直接依赖它的 test_m65e_kernel_timeline.py 未能导入。

性能收尾 artifacts：

build-perf/qwen3-8b-tp2/final-b1-five-round-summary.json，SHA256 7a5505e7edb348048f2075786fa73eac672c113e0391a52470e8ba554894c053；
build-perf/qwen3-8b-tp2/final-b4-five-round-summary.json，SHA256 71483ffc25023d4f21f9703ac20556a0ce5b4149301158242ac588393d5dcbd4；
build-perf/qwen3-4b-tp2/final-b1-five-round-summary.json，SHA256 69d7da0f9d8ce2889a71460b958ca0c64bd56fe4b67a0caf8cabdecfb66ad609；
build-perf/qwen3-8b-tp2/decode-gemm-quad-candidate.json，SHA256 8314eb3c5da212c44029526418186863fc023c7e9a8fa0fbf211727c3bea68a4；
build-perf/qwen3-8b-tp2/timeline-final/results_final.csv，SHA256 0f4cfda0e2bc1ff356baf87a02955c9ffe0d47ee5052a31b1279b19970edc4e0；
build-perf/m65f-topk-tail-benchmark.json，SHA256 a24f324f2b1ef64a26de5151de624c47d1b047df63d057f6f82bbd88bc96a45a。

Qwen3 TP4 与 RCCL

TP4 将原来固定两个 Rank 的执行器泛化为动态 world_size。服务仍是单进程，但为四个 HIP device 分别创建 Backend、Stream、权重 shard、Runner、Paged KV Cache 和 InferenceEngine；TensorParallelInferenceEngine 并发推进四个 Rank，并在任一 Rank 失败时对整个通信组执行 abort。TP2 继续使用 HIP P2P event/kernel collective，TP4 使用 DTK 自带的 RCCL 2.13.4+hip5.7（/opt/dtk/lib/librccl.so）。

Transformer 层保持标准 Megatron TP 语义：QKV 和 Gate/Up 做列切分，O 和 Down 做行切分，行切分结果通过 BF16 AllReduce 汇总。LM Head 按 vocabulary 行切到四个 Rank，每个 Rank 计算局部 F32 logits，再用 AllGather 形成 batch-row-major 完整 logits。RCCL 实现另外覆盖 F32/BF16 AllReduce、F32 二维 AllGather、Int32 Broadcast、融合 reduce-add/reduce-add-RMSNorm 和 group-wide abort。独立四卡 smoke 与 framework test 均通过；后者还验证了 descriptor rendezvous 和 rank-major 到 batch-row-major 重排。

HTTP/OpenAI 接口无需增加 TP4 专用协议，通过 --devices 0,1,2,3 启动后仍提供 /health、/v1/models、/v1/completions、/v1/chat/completions、SSE 和 metrics。

Qwen3-8B TP4 基线

8B TP4 服务已验证普通 completion 和 SSE chat，首 token 保持为 ID 151667 （<think>）。服务加载后四卡显存约为 5.65-5.75 GiB/卡，退出后均回到约 2 MiB。 hidden=4096、20 次 warmup、200 次计时的 RCCL AllReduce microbenchmark：

Workload	Payload	Latency
Decode B1	8 KiB	47.59 us
Decode B4	32 KiB	120.87 us
Prefill 32	256 KiB	136.30 us
Prefill 256	2 MiB	340.56 us

同一 persistent server 上四种场景各运行五轮，TPOT 使用首 token 后生成窗口的平均值：

场景	TTFT p50/p95	TPOT p50/p95	Completion throughput
B1, 32 words, 8 output	349.586 / 354.013 ms	8.320 / 8.372 ms	19.616 token/s
B1, 128 words, 8 output	549.064 / 550.912 ms	8.805 / 8.831 ms	13.100 token/s
B4, 32 words, 16 output	590.010 / 598.375 ms	23.723 / 23.898 ms	67.596 token/s
B4, 128 words, 8 output	1155.557 / 1201.775 ms	10.694 / 11.158 ms	25.373 token/s

正式 artifact 为 build-perf/qwen3-8b-tp4/five-round-summary.json，SHA256 f35120f39b9ac7f69aede8627b72016276dbd026d290f88ffd019c192df057fa。

Qwen3-14B TP4 Bring-up

14B 使用官方 revision 40c069824f4251a91eefaf281ebe4c544efd3e18，严格 profile 为 hidden_size=5120、intermediate_size=17408、40 layers、40 Query Heads、8 KV Heads、head_dim=128、vocab_size=151936，且 Embedding 与 LM Head 不共享权重。八个 Safetensors shard 共 29,536,665,640 bytes。TP4 Rank-local shape 为 10 Query Heads、2 KV Heads、intermediate_size=4352 和 vocabulary_size=37984。manifest 固定 13 个文件的大小与 SHA256，其自身 SHA256 为 2d59d2caaf53c4206159d7980b66d62db258f19cfd407efa37d21649581e9141。

正式服务使用下面的受控容量，避免默认 64 * 4096 KV 容量占用大量显存：

build-perf/bin/metainfer_ref_qwen3_server \
  --model-dir "$HOME/.cache/metainfer/models/Qwen3-14B" \
  --devices 0,1,2,3 --host 127.0.0.1 --port 18085 \
  --max-context-tokens 256 --max-output-tokens 16 \
  --max-active-sequences 4 --max-queue-size 16 \
  --max-batched-tokens 256 --prefill-chunk-tokens 128 \
  --block-size-tokens 16

端到端验证中，/health 和 /v1/models 正确报告 Qwen3-14B；非流式 prompt Hello 生成 : I have a，SSE chat 逐帧生成 <think>\nOkay, 并以 [DONE] 结束。单 active 配置加载后四卡约 9.00-9.03 GiB/卡，四 active 正式配置峰值约 9.09-9.11 GiB/卡；SIGINT 正常退出后四卡均回到约 2 MiB。

hidden=5120 的同方法 RCCL microbenchmark：

Workload	Payload	Latency	Payload bandwidth
Decode B1	10 KiB	36.28 us	0.282 GB/s
Decode B4	40 KiB	107.95 us	0.379 GB/s
Prefill 32	320 KiB	135.73 us	2.414 GB/s
Prefill 256	2.5 MiB	388.44 us	6.749 GB/s

五轮正式 HTTP/OpenAI 基线：

场景	TTFT p50/p95	TPOT p50/p95	Completion throughput
B1, 32 words, 8 output	510.646 / 540.426 ms	7.753 / 9.192 ms	14.159 token/s
B1, 128 words, 8 output	824.300 / 825.146 ms	8.299 / 8.326 ms	9.065 token/s
B4, 32 words, 16 output	898.310 / 900.419 ms	26.814 / 27.101 ms	49.198 token/s
B4, 128 words, 8 output	1840.096 / 1889.297 ms	11.592 / 12.199 ms	16.415 token/s

相对同协议的 8B TP4，14B 的短/长 B1 TTFT 分别增加 46.1%/50.1%，B4 decode 和 mixed TTFT 增加 52.3%/59.2%；这主要来自 40 层、更大的 Hidden/MLP 权重以及更大的 prefill GEMM。B4 decode TPOT 增加 13.0%，吞吐降低 27.2%；B4 mixed TPOT 增加 8.4%，吞吐降低 35.3%。短 B1 的观测 TPOT 反而低 6.8%，但整个请求吞吐仍低 27.8%，说明该短生成窗口主要受 TTFT、服务固定开销和测量边界影响，不能据此判断 14B 的 Decode Kernel 比 8B 更快。

Qwen3-14B Linear Shape 优化

初始生产策略中的 medium-M 白名单只覆盖 4B/8B，14B Prefill 因此全部回退 hipblasGemmEx；Decode 虽已使用 BF16x2 small-M，但 BF16x4 白名单同样没有 14B Rank-local shape。新增 benchmark profile 精确覆盖 QKV [5120 -> 1792]、O [1280 -> 5120]、Gate/Up [5120 -> 8704] 和 Down [4352 -> 5120]，所有候选均与 hipBLAS BF16 输出一致且 steady-state allocation 不变。

Decode 只选择有明确余量的 BF16x4 case：B1 QKV/O/Down 相对 BF16x2 分别为 1.28x/1.13x/1.07x，B4 只选择 O 的 1.59x。B1 Gate/Up 只有 1.01x，B4 QKV、 Gate/Up、Down 分别为 0.87x/0.86x/0.88x，继续使用 BF16x2。

Prefill 对 M=5/8/9/16/24/32/40/48/64/128 测量 group-8/16/32。生产路由固定为：

M=5..9：四类投影使用 group-8；
M=10..32：QKV、Gate/Up 使用 group-16，O、Down 使用 group-8；
M=33..48：QKV 使用 group-16，O、Down 使用 group-8，Gate/Up 回退 hipBLAS；
M>=49：全部回退 hipBLAS。

入选候选在 M=5 的加速范围为 5.24x-8.18x，M=16 为 2.17x-3.47x，M=32 为 1.10x-1.85x，M=40 为 1.30x-1.53x，M=48 为 1.11x-1.32x。M=64/128 的 medium-M 已持平或明显退化，因此没有外推白名单。

相同 persistent server、相同四场景的五轮复测：

场景	优化后 TTFT p50/p95	优化后 TPOT p50/p95	优化后吞吐	相对基线
B1, 32 words, 8 output	449.651 / 480.410 ms	7.537 / 7.560 ms	15.930 token/s	TTFT -11.9%，TPOT -2.8%
B1, 128 words, 8 output	531.088 / 532.021 ms	8.091 / 8.130 ms	13.614 token/s	TTFT -35.6%，吞吐 +50.2%
B4, 32 words, 16 output	833.674 / 842.310 ms	25.973 / 26.600 ms	52.290 token/s	TTFT -7.2%，吞吐 +6.3%
B4, 128 words, 8 output	1546.593 / 1592.620 ms	11.397 / 11.814 ms	19.350 token/s	TTFT -16.0%，吞吐 +17.9%

四个场景均无回退，SSE 固定请求继续逐 token 生成 <think>\nOkay,。候选 artifacts：

build-perf/qwen3-14b-tp4/decode-linear-candidates.json，SHA256 92baf3b21d07869484d0b247529199f55721f10a169dda9f473b0ca80f7581bc；
build-perf/qwen3-14b-tp4/prefill-linear-candidates.json，SHA256 83b95f0dddbb2cbba50d6db4411864d6cb60307129a599690a080988fb815d98；
build-perf/qwen3-14b-tp4/linear-policy-five-round-summary.json，SHA256 78f79cbbe8f016553ec7caf910bc37b155aad9a07f4ac9d5ed79a7cc1baf241b。

14B artifacts：

build-perf/qwen3-14b-tp4/five-round-summary.json，SHA256 90fec01bd1880d51aa65ba2c0ee6a6b79bdd9f49621e09c0a13c95bc674f4c62；
build-perf/qwen3-14b-tp4/collective-benchmark.json，SHA256 3c9714a2fd67d39024408f2f14f793f52c6eeb2ebd500d15843a68ddc5419fed。

Qwen3-14B TP4 双 RCCL communicator

新增的 Rank-specific timeline driver 在模型加载和 warmup 后，分别采集 B1/B4、短/长 Prefill 和 Decode 的八个固定窗口，并按 rocprof GPU_ID 独立聚合四个 Rank。初始 Decode timeline 中，B1 的 GEMM、collective 占 Rank 0 kernel time 的 61.4%/17.7%；B4 L32 为 50.9%/36.5%，B4 L128 为 42.2%/43.7%。B4 每步 80 次 40 KiB AllReduce 累计 15.48-23.10 ms，而 LM Head AllGather 只有 0.22-0.41 ms，因此这一阶段没有优先实现收益上限不足 1 ms 的 distributed argmax。

本机 RCCL 2.13.4 的默认调优模型会让 40 KiB 选择高带宽、高启动延迟协议。全局强制 NCCL_PROTO=LL/LL128 虽能把 40 KiB 从约 110 us 降到约 39 us，但会把 2.5 MiB 从约 389 us 恶化到约 946 us。生产实现因此创建两个 RCCL communicator：默认 communicator 处理大消息，初始化时固定为 LL128 的 communicator 只处理不超过 64 KiB 的 AllReduce。创建期间的 NCCL_PROTO 修改受进程内 mutex 保护，并在工厂返回前恢复，不会改变服务或调用方环境；AllGather 和 Broadcast 继续使用默认 communicator。

hidden=5120、20 次 warmup/200 次测量的 microbenchmark：

Payload	原延迟	双 communicator	变化
10 KiB	36.28 us	28.80 us	-20.6%
40 KiB	107.95 us	37.22 us	-65.5%
320 KiB	135.73 us	135.39 us	-0.2%
2.5 MiB	388.44 us	388.40 us	持平

优化后 Rank 0 collective kernel time 在 B1 L32/L128 下降 22.9%/15.1%，B4 L32/L128 下降 25.0%/37.4%；B4 L128 的受控 step wall time 下降 10.6%。相同 persistent server 和五轮 HTTP/OpenAI workload 的正式结果：

场景	TTFT p50/p95	TPOT p50/p95	吞吐	相对 Linear Shape 版本
B1, 32 words, 8 output	444.699 / 477.602 ms	7.223 / 7.265 ms	16.155 token/s	TTFT -1.1%，TPOT -4.2%
B1, 128 words, 8 output	526.067 / 531.639 ms	7.798 / 7.837 ms	13.773 token/s	TTFT -0.9%，TPOT -3.6%
B4, 32 words, 16 output	801.191 / 818.862 ms	21.726 / 22.634 ms	56.698 token/s	TPOT -16.4%，吞吐 +8.4%
B4, 128 words, 8 output	1510.130 / 1554.001 ms	9.401 / 10.009 ms	19.971 token/s	TPOT -17.5%，吞吐 +3.2%

固定 SSE 请求继续逐 token 生成 <think>\nOkay,，服务退出后四卡显存均回到约 2 MiB。新增 artifacts：

build-perf/qwen3-14b-tp4/kernel-timeline-baseline.json，SHA256 e9147419b926e3eeaa66b68819a9f5cfcba65eadc98964c1a0c2930f7166bf4c；
build-perf/qwen3-14b-tp4/kernel-timeline-dual-rccl.json，SHA256 9417c5d1ca45557f14f9d10e71f83cc464db47761da68473c1bfec82f6d6909b；
build-perf/qwen3-14b-tp4/dual-rccl-five-round-summary.json，SHA256 04728a103b31c8fae65bcdb5c0b5200c35d0ff1c5e367c7edf5d30a41c874518。

本阶段最终验证包括通用 CTest 136/136、TP4 RCCL 2/2、TP2 collective 9/9、 HIP Backend 10/10、HIP Operators 21/21、Dynamic Engine HIP 10/10 和 Tiny HIP 6/6，所有测试均通过。当前 TP4 的主要工程限制是单进程四 Rank，尚无多进程启动、通信与计算重叠、量化权重或跨节点执行。

Qwen3-14B TP4 B4 multi-output-row GEMV

双 RCCL communicator 将通信成本压低后，B4 Decode 的主要计算热点变为 PackedSmallMBf16LinearKernel<4>。原 kernel 让每个 wave 计算一个输出通道；新的 PackedRowsB4Bf16LinearKernel 让一个 wave 同时计算 2、3 或 4 个输出通道，从而让这些输出通道复用同一份 B4 input load。每个输出通道内部仍保持原 BF16x2 load、FMA 和 wave reduction 顺序，因此与 hipBLAS 的 BF16 输出逐元素一致，steady-state allocation 也没有增加。

对 14B TP4 rank-local 四种投影执行 20 次 warmup、200 次测量，共重复五轮。生产路由只固化有稳定收益的精确 B4 shape，不向相近维度外推：

投影	原生产 kernel 中位数	新生产 kernel 中位数	五轮收益范围	路由
QKV `[5120 -> 1792]`	0.05805 ms	候选均退化	-	保留 BF16x2
O `[1280 -> 5120]`	0.03769 ms	0.03598 ms	`1.040x-1.091x`	rows/wave=3
Gate/Up `[5120 -> 8704]`	0.24274 ms	0.21842 ms	`1.107x-1.114x`	rows/wave=2
Down `[4352 -> 5120]`	0.11934 ms	0.09053 ms	`1.315x-1.329x`	rows/wave=3

新 Kernel Timeline 中，B4 L32 每个 Rank 均出现 40 次 rows=2、80 次 rows=3，原 BF16x2 只剩 40 次 QKV，说明 O/Gate-Up/Down 路由均已生效。Rank 0 的四类 Decode linear kernel 总时间从 18.736 ms 降到 16.076 ms，下降 14.2%。这一轮 profile 中的 RCCL 和 host gap 波动较大，单次 range wall time 从 82.109 ms 增至 88.919 ms，因此不把单次 profile wall time 当作端到端结论，而以 persistent server 五轮结果裁决。

相同服务配置、相同四场景的正式五轮结果如下：

场景	TTFT p50/p95	TPOT p50/p95	吞吐	相对双 RCCL 版本
B1, 32 words, 8 output	444.300 / 446.872 ms	7.235 / 7.289 ms	16.154 token/s	p50 基本持平
B1, 128 words, 8 output	526.942 / 527.874 ms	7.804 / 7.880 ms	13.746 token/s	p50 基本持平
B4, 32 words, 16 output	786.342 / 794.883 ms	19.803 / 19.901 ms	59.066 token/s	TTFT -1.9%，TPOT -8.9%，吞吐 +4.2%
B4, 128 words, 8 output	1498.562 / 1554.846 ms	8.969 / 9.193 ms	20.155 token/s	TTFT -0.8%，TPOT -4.6%，吞吐 +0.9%

固定 SSE 请求仍逐 token 输出 <think>\nOkay,。四 active 服务运行时四卡使用 9.00-9.05 GiB/卡，SIGINT 退出后均回到约 2 MiB。新增 artifacts：

build-perf/qwen3-14b-tp4/multi-row-round-{1..5}.json；
build-perf/qwen3-14b-tp4/kernel-timeline-packed-rows.json，SHA256 5a104b029ffdb1c72ceb0efd53b30868e2496213a11dc9ddf34f131285a9a280；
build-perf/qwen3-14b-tp4/packed-rows-five-round-summary.json，SHA256 ebb510396dcd84ef1c9fa6f2473f60cfe632982f0530076520610c0a91d7bb32。

本阶段回归包括完整 HIP 构建、通用 CTest 136/136、HIP Backend 10/10、HIP Operators 21/21、Dynamic Engine HIP 10/10、Tiny HIP 6/6、TP2 collective 9/9、TP4 RCCL 2/2 和 timeline tooling 2/2，全部通过。

Qwen3-14B TP4 Persistent Rank Workers

原 TensorParallelInferenceEngine::Tick() 每轮都会创建 world_size-1 个 peer Rank 线程，完成后立即 join；TP4 的每个 Prefill/Decode Tick 因而重复创建和销毁三个线程。现在 Engine 在创建时启动 Rank 1-3 的常驻 worker，通过 epoch 和 condition variable 下发 Tick；Rank 0 仍在服务调用线程执行。每 Rank 结果槽同样跨 Tick 复用，不再为结果 vector 做每轮分配。

新的生命周期保持原错误语义：任一 peer 失败仍先执行 group-wide collective abort，主线程等待所有 peer 发布结果后再校验 status、stats 和 events；Shutdown 或析构都会唤醒并 join 空闲 worker。Rank 0 或 peer 的 Tick 异常统一转换为带 Rank 信息的 Internal，不会逃出调用线程或在 peer 线程触发 std::terminate。TDD 测试使用 peer callback 的 thread-local 初始化计数，旧实现连续两次 Tick 得到 2，新实现得到 1，直接证明复用的是线程生命周期而不是偶然相同的 thread ID。

现有 Rank-specific Kernel Timeline driver 直接调用四个 Runner，并拥有自己的一组临时 Rank 线程，不经过 TensorParallelInferenceEngine。因此它适合继续测 GPU kernel，但不能用于归因本阶段的 Engine host 调度收益。本阶段生产裁决以同一 persistent HTTP 服务的五轮 A/B 为准：

场景	TTFT p50/p95	TPOT p50/p95	吞吐	相对 multi-output-row 版本
B1, 32 words, 8 output	442.145 / 445.833 ms	7.139 / 7.149 ms	16.258 token/s	TTFT -0.5%，TPOT -1.3%，吞吐 +0.6%
B1, 128 words, 8 output	524.203 / 538.147 ms	7.715 / 7.750 ms	13.830 token/s	TTFT -0.5%，TPOT -1.1%，吞吐 +0.6%
B4, 32 words, 16 output	783.202 / 785.569 ms	19.594 / 19.635 ms	59.354 token/s	TTFT -0.4%，TPOT -1.1%，吞吐 +0.5%
B4, 128 words, 8 output	1496.163 / 1543.690 ms	8.867 / 8.998 ms	20.191 token/s	TTFT -0.2%，TPOT -1.1%，吞吐 +0.2%

B4 Decode 每轮 TPOT 为 19.504-19.617 ms，上一版本为 19.709-19.862 ms，五轮区间没有重叠；最终四场景的 TTFT、TPOT 和吞吐均无回退。因此保留 Persistent Rank Worker，但把它定位为约 1% 的调度优化，不把 profile 中全部 host gap 归功于线程复用。固定 SSE 仍输出 <think>\nOkay,，退出后四卡均回到约 2 MiB。

正式 artifact 为 build-perf/qwen3-14b-tp4/persistent-rank-workers-five-round-summary.json，SHA256 94a844835266cead33dd4ea88232473c1ad7cbef259fdd39816e95e74115b640。完整 HIP 构建、通用 CTest 139/139、Tensor Parallel Engine 8/8、HIP Backend 10/10、HIP Operators 21/21、Dynamic Engine HIP 10/10、Tiny HIP 6/6、TP2 collective 9/9、TP4 RCCL 2/2 和 timeline tooling 2/2 均通过。

Qwen3-14B TP4 B4 tiled LDS input cache

PackedRowsB4Bf16LinearKernel 中的四个 wave 原本会各自从全局显存重复读取同一份 B4 输入。新的 PackedRowsB4LdsBf16LinearKernel 让 256-thread block 先协作加载一块输入到 LDS，再由四个 wave 共享。候选保持每个输出通道原有的 BF16x2、FMA 与 wave reduction 顺序，并支持 strided B4 input。最后一个输出块中的无效 wave 仍参与所有 barrier；最后一个 input tile 之后省略不再需要的 barrier。

实验接口覆盖 rows-per-wave=2/3 与 tile-pairs=256/512/1024。对四种 14B TP4 rank-local Decode shape 执行 20 次 warmup、200 次测量并重复五轮后，只有 Gate/Up [B4, 5120 -> 8704] 稳定超过 3% 生产门槛：rows-per-wave=2、tile-pairs=256 的中位耗时为 0.188140 ms，相对原 multi-output-row kernel 的五轮加速范围为 1.1611x-1.1650x。O/Down 的收益最多约 1.7%，QKV 退化，因此三者不进入 LDS 生产路由；生产策略只对白名单中的精确 Gate/Up shape 生效。后续 benchmark schema 6 把比较基准明确记录为 pre-lds-production，并为每个候选输出 baseline_kernel、baseline_ms 和 vs_baseline，避免把当前生产路由误标成旧基线。六种候选在五轮中均通过 hipBLAS BF16 correctness 检查，且每个 case 的 allocation_stable=true。

Rank-specific Kernel Timeline 中，B4 L32 的每个 Rank 都出现 40 次 PackedRowsB4LdsBf16LinearKernel<2,256>。Rank 0 Gate/Up kernel 总时间从 8.579 ms 降到 7.763 ms（-9.5%），四类 small-M linear 合计从 16.076 ms 降到 15.287 ms（-4.9%）。RCCL 与 host gap 在单次 profile 中仍有波动，因此生产裁决继续使用 persistent HTTP 服务的五轮 A/B：

场景	TTFT p50 变化	TPOT p50 变化	Completion throughput 变化
B4, 32 words, 16 output	`-0.56%`	`-2.70%`	`+1.18%`
B4, 128 words, 8 output	`-0.18%`	`-1.46%`	`+0.29%`
B1, 32 words, 8 output	`+0.10%`	`+0.43%`	`-0.21%`
B1, 128 words, 8 output	`+0.09%`	`+0.11%`	`-0.04%`

B1 不命中 LDS 白名单，其变化视为测量噪声。固定 SSE 请求继续输出 <think>\nOkay,。本阶段 artifacts：

build-perf/qwen3-14b-tp4/lds-refined-round-{1..5}.json，五轮 SHA256 依次为 94b124fc348c4140281d21913cb22967fa6bb7b7b0d0dcbf8b0baf8226ba3e08、 67e73f48e46a5877a47c0a4549e318e470b8fb80aa69bd945405ae2207f87d8f、 2a9307cd99f9f78a79e9317f79e109b2206633b774548b1f7a892f8d3964ff92、 ff197077f44c1a3cb40677c4ca9378952d8541fde399ca91c0f87408791cecf8、 48035ddbe72b996dd5f0e9624089f57a8dc0a1fa7ede32d9e11e4f8727d62433；
build-perf/qwen3-14b-tp4/kernel-timeline-lds-gate.json，SHA256 2fca8b236dce85835ddea6e57d575696d9eca31e0023efb219291e4adf2ad49a；
build-perf/qwen3-14b-tp4/lds-gate-five-round-summary.json，SHA256 f0168e146c5a6e53b41ae7fbf42c24e067271ace2515c7b7ccec710b18766912。

最终回归包括完整 HIP 构建、通用 CTest 139/139、HIP Backend 10/10、HIP Operators 21/21、Dynamic Engine HIP 10/10、Tiny HIP 6/6、TP collective 9/9、TP4 RCCL 2/2 和 timeline tooling 2/2，全部通过。

Qwen3-14B TP4 BF16x2 Paired Decode Attention

14B TP4 的 rank-local Decode Attention 是精确的 10Q/2KV/head-dim=128。原 fused kernel 每个 lane 分别读取第 lane 和 lane+64 列；新 paired kernel 保留一 block 对应一个 Query Head、在线 Softmax、BF16 score rounding 和 Paged KV 布局，但让每个 lane 通过对齐的 __hip_bfloat162 连续读取 2*lane、2*lane+1。实验比较每个 head 使用 1/2/4/8 个 wave，W8 使用 512 threads，并以局部 __launch_bounds__(512) 保证编译和启动约束一致。空 wave 仍写入中性统计量并参与统一 barrier/merge，不引入 score、 probability 或 packed workspace。

生产路由只对 query_heads=10、kv_heads=2、head_dim=128、偶数 query row stride，以及 query/KV pool/output 全部 4-byte 对齐的输入选择 W8；其他 shape 或 alignment 全部回退原 scalar fused kernel。Trusted API 仍要求调用方保证每个被 past_length + 1 覆盖的 logical block table entry 有效且落在 physical block pool 范围内；未使用的 table entry 可以为 -1，kernel 不应读取它们。

精确 shape microbenchmark 对 B1/B4 和 context 32/128/256 分别执行 20 次 warmup、 200 次测量并重复五轮。所有候选均通过 BF16 correctness 和 steady-state allocation 检查；生产门槛要求每个 case 至少 0.97x，并以 batch_size * context_tokens 为权重计算几何均值。W8 五轮加权几何均值为 1.4994x-1.5026x，是唯一稳定超过 1.03x 门槛的候选。逐场景五轮中位数如下：

场景	Scalar fused	Paired W8	加速
B1 L32	0.026389 ms	0.019907 ms	1.3260x
B1 L128	0.062702 ms	0.042405 ms	1.4781x
B1 L256	0.112772 ms	0.073028 ms	1.5437x
B4 L32	0.025130 ms	0.019881 ms	1.2637x
B4 L128	0.062507 ms	0.042381 ms	1.4742x
B4 L256	0.112821 ms	0.072829 ms	1.5486x

Rank-specific Kernel Timeline 的四个 Decode range、四个 Rank 均出现 40 次 FusedPagedDecodeAttentionPairsKernel<8>，Prefill range 不出现该 kernel。相对上一版 LDS Gate，Rank 0 的 paired attention kernel 与完整 attention family 变化为：

Range	Attention kernel	Attention family
B1 L32	-26.7%	-21.0%
B1 L128	-38.4%	-35.4%
B4 L32	-27.9%	-22.5%
B4 L128	-40.1%	-36.9%

同一 persistent TP4 HTTP 服务的正式五轮 A/B 继续以 LDS Gate 为基线。四个场景的 p50 与 completion throughput 均改善：

场景	TTFT p50	TPOT p50	Completion throughput
B4 Decode	-0.40%	-1.60%	+0.66%
B4 Mixed	-0.72%	-4.08%	+0.82%
B1 Prefill	-1.24%	-4.21%	+1.53%
B1 Short	-0.16%	-1.37%	+0.35%

固定 SSE 输出仍以 <think>\nOkay, 开头；服务退出后四卡均回到约 2 MiB。正式 artifacts 为：

build-perf/qwen3-14b-tp4/paired-attention-round-{1..5}.json，SHA256 依次为 3e8c8742e4093f6630a19cb670acea11243b9c03a4fefd5ab6dc25cb29759a07、 0f48fe2bb82cbc16c114164011c437e4bb3c39d9367e64dbe79629882b8b15bd、 2550364f21454e5819142f6a132a997f36e442e66772011ae0329a9431a857ad、 ef844938456b203342368823cdf408bba6bda2640da499ec1d4df639932d4b0f、 fea443935fe50c1d29787ea332a944324449e8730b0ab6642950057c35f4b55f；
build-perf/qwen3-14b-tp4/kernel-timeline-paired-attention.json，SHA256 ddaa0c2681a0d64d75e6d3c61e8ede9b39f25983e57f89f6094a32b247e911eb；
build-perf/qwen3-14b-tp4/paired-attention-five-round-summary.json，SHA256 d7f13db2283b7a2690b7151f27d7a3574ada114075d9066962e43666ad5acfb7。

最终回归包括完整 HIP 构建、通用 CTest 139/139、HIP Backend 10/10、HIP Operators 22/22、Dynamic Engine HIP 10/10、Tiny HIP 6/6、TP collective 9/9、TP4 RCCL 2/2，以及 benchmark/Timeline tooling 3/3，全部通过。

Qwen3-14B TP4 Final BF16 Decode 收尾

最后一轮按 Decode GEMM/GEMV、RCCL 小消息和 Elementwise launch fusion 三条路径逐项裁决。Decode 新增 128-bit BF16x8 OctetSmallMBf16LinearKernel，五轮 20 warmup / 200 iterations 只选择精确 B1 QKV [5120 -> 1792] 和 Gate/Up [5120 -> 8704]：相对原生产 kernel 的收益范围分别为 1.0740x-1.0784x 和 1.0617x-1.0841x。O/Down 没有在五轮中稳定超过 3%，继续使用 BF16x4。

RCCL 新增 ordered trusted 实验入口，并让 fused reduce-add 直接以 branch 为 send buffer、预分配 scratch 为 receive buffer，消除原 D2D copy。五轮结果证明跳过 Host descriptor rendezvous 的收益不足以切换生产入口：

Payload	Validated 中位数	Trusted 中位数	Trusted 加速
10 KiB	28.7666 us	28.4790 us	1.0101x
40 KiB	37.9250 us	37.1612 us	1.0206x
320 KiB	135.5655 us	135.4174 us	1.0011x
2.5 MiB	388.3360 us	387.9352 us	1.0010x

因此 trusted 保留为实验 API，Runner 继续走 validated 路径；生产 RCCL 仍使用既有的双 communicator 和 64 KiB 以下 LL128 fast path。

Elementwise 路径新增 hidden=5120、B1-B4 精确 shape 的 BF16x2 paired AddRmsNorm。五轮加权几何平均加速为 1.7116x-1.7209x，updated residual 逐位一致、normalized 输出零误差且 allocation delta 为 0。随后把 normalization 跨层流水化：Embedding 后只做一次 layer-0 Norm，每层 attention tail 和 MLP tail 各执行一次 paired AddRmsNorm，后一项同时产生下一层 normalized hidden；最后一层直接使用 final norm weight。

最终 Rank-specific Timeline 在四个 Decode range、四个 Rank 上都满足：

B1 每步 80 次 BF16x8 linear，B4 不命中该白名单；
每步 80 次 AddRmsNormPairsKernel；
AddKernel 从 40 次降到 0，NormKernel 从 41 次降到 1；
总 kernel launch 从 570 降到 530；
Rank 0 B1 elementwise family 下降 33.6%-34.4%。

相对 paired-attention 版本，同一 persistent TP4 HTTP 服务五轮结果为：

场景	最终 TTFT p50/p95	最终 TPOT p50/p95	最终吞吐	p50/吞吐变化
B1, 32 words, 8 output	433.017 / 465.942 ms	6.638 / 6.663 ms	16.678 token/s	TTFT -2.00%，TPOT -6.13%，吞吐 +2.45%
B1, 128 words, 8 output	510.034 / 520.567 ms	6.922 / 6.971 ms	14.316 token/s	TTFT -1.57%，TPOT -6.44%，吞吐 +1.99%
B4, 32 words, 16 output	767.585 / 776.526 ms	17.891 / 17.932 ms	61.656 token/s	TTFT -1.05%，TPOT -4.63%，吞吐 +1.99%
B4, 128 words, 8 output	1473.912 / 1511.904 ms	7.996 / 8.038 ms	20.600 token/s	TTFT -0.60%，TPOT -4.60%，吞吐 +0.90%

B1 short 的 p95 来自首轮 465.942 ms 冷启动离群值，其余四轮为 432.837-433.416 ms。固定 SSE 仍以 <think>\nOkay, 开头，服务退出后四卡均回到约 2 MiB。正式 artifacts：

build-perf/qwen3-14b-tp4/final-bf16-gemm-round-{1..5}.json；
build-perf/qwen3-14b-tp4/final-elementwise-round-{1..5}.json；
build-perf/qwen3-14b-tp4/final-rccl-trusted-five-round-summary.json，SHA256 3e49d8bc2c667bc070fad95990c3a588389bed5f3a64c78d5c9995521edaf92e；
build-perf/qwen3-14b-tp4/kernel-timeline-final-bf16.json，SHA256 e939001cb27bf5a28a524efeaacf0711ec125dd4957d065eb0f74e37eee1cc13；
build-perf/qwen3-14b-tp4/final-bf16-five-round-summary.json，SHA256 93f73b48d2cae91ab06e4a3c4f1ccd58ad6e37794251465575f3cd3356822d32。

最终验证包括完整 HIP build、通用 CTest 139/139、HIP Backend 11/11、HIP Operators 23/23、Dynamic Engine HIP 10/10、Tiny HIP 6/6、TP collective 9/9、TP4 RCCL 2/2，以及 collective/elementwise/Timeline tooling，全部通过。

Qwen3-14B TP4 16x4096 KV Capacity

为区分 KV 压缩与服务容量削减，正式测试了 16 active × 4096 context。TP4 每 Rank 有 2 个 KV Head，BF16 KV 每 token/卡为 40 KiB，因此 65,536 个 token slots 对应 2.5 GiB/卡，物理池为 4096 个 16-token block。服务使用 256-token Chunked Prefill、 256 output 上限和 32 HTTP connection limit。

max_batched_tokens=2048 配置成功启动，空载显存为 11883-11931 MiB/卡，压力期间峰值为 11920-11969 MiB/卡。相同 B1/B4 常规 workload 五轮相对 4x256 最终基线基本持平：B1 short TTFT -0.06%、TPOT +0.07%，B4 Decode TTFT -0.03%、 TPOT +0.47%；其他 TTFT/吞吐改善属于运行波动，没有发现因 KV Pool 扩大导致的稳定性能回退。

容量压力使用 16 个并发请求，每请求 4080 个 runtime words；实际 tokenizer 输出 4088 prompt tokens 和 1 completion token。总 prompt 为 65,408 tokens，占物理容量 99.8047%。16/16 请求均返回 HTTP 200，未发生 OOM、RCCL abort 或 KV exhaustion。

同时比较两种 Prefill 调度上限：

max_batched_tokens	空载/峰值 VRAM	总墙钟	Prompt 吞吐	请求延迟 min/median/max
2048	11.60-11.65 / 11.64-11.69 GiB	214.006 s	305.636 token/s	99.085 / 156.929 / 213.995 s
4096	11.71-11.76 / 11.74-11.79 GiB	211.306 s	309.541 token/s	194.455 / 211.294 / 211.301 s

4096 把 16 个 256-token chunk 放进同一个公平推进组，总吞吐只提高 1.28%、尾延迟降低 1.26%，但请求中位延迟恶化 34.64%；2048 会分成两组八请求，使首组约 99-108 秒完成。因此当前默认推荐 max_batched_tokens=2048：容量相同、总吞吐接近，但能更早返回一半请求。正式 artifacts 为：

build-perf/qwen3-14b-tp4/kv-capacity-16x4096-five-round-summary.json，SHA256 52481d6de910f5a149a58521e181b48ff2fcb863d52e66936fb4c1d423832d68；
build-perf/qwen3-14b-tp4/kv-capacity-16x4096-stress-summary.json，SHA256 39f0e5067678996164b434e12c4ad4fdbf6276e22565a432935077612a397ae8。

Qwen3-14B TP4 Performance Matrix

在上述 16 active x 4096 context 容量配置上，使用同一个 persistent HTTP 服务补充上下文长度、并发、排队、长短混合、稳定性、输出长度、动态到达和 Prefill chunk 八类实验。除 chunk sweep 外，服务固定为 max_batched_tokens=2048、 prefill_chunk_tokens=256；TTFT 从请求提交到第一个非空 SSE token，TPOT 为首 token 之后服务端生成窗口的平均值。Prompt 吞吐仍以重复英文 word 数作为稳定工作量代理，不能与 tokenizer token/s 直接等同。

单请求 Context 扫描，每档两轮取中位数：

Prompt words	TTFT	TPOT	Prompt words/s
32	449.606 ms	6.631 ms	64.498
128	494.842 ms	6.958 ms	235.238
512	1415.150 ms	8.216 ms	347.516
1024	2855.955 ms	9.931 ms	349.945
2048	6301.020 ms	13.248 ms	320.276
4080	15253.820 ms	19.822 ms	265.047

512-1024 words 达到约 348-350 words/s 的 Prefill 甜点；从 2048 增长到 4080 words 时 TTFT 增至 2.42x，已经出现 Attention 随上下文增长的非线性成本。

固定 128-word prompt、32-token output 的并发扫描，每档三轮取中位数：

并发	TTFT	TPOT	Completion token/s	Prompt words/s
1	494.524 ms	20.471 ms	28.324	113.298
4	1396.279 ms	23.716 ms	59.997	239.989
8	2701.471 ms	67.071 ms	53.501	214.006
16	5066.536 ms	119.721 ms	58.273	233.092

B4 相对 B1 的 completion throughput 提高 111.8%，是当前 Decode 吞吐甜点；B8 和 B16 分别比 B4 低 10.8% 和 2.9%，同时 TTFT/TPOT 大幅上升。32 个并发请求压入 16 个 active slot 的两轮排队实验均 32/32 完成，总吞吐中位数为 39.315 token/s，TTFT p50/p95 约为 8.256/11.586 s，证明队列容量有效，但过载时尾延迟会显著增长。B4、128-word、8-output 连续十轮吞吐范围为 21.981-22.028 token/s，变异系数仅 0.0674%，稳态重复性良好。

固定 B1、128-word prompt 的输出长度扫描显示，8/32/128 output 的 TTFT 保持在 494.441/494.578/497.923 ms，说明输出上限不改变 Prefill 固定成本；completion throughput 随固定成本被摊薄，从 14.716 提高到 28.129/36.022 token/s，128-token 长 Decode 的稳态平均 TPOT 为 24.019 ms。2-token 档用于首 token 边界检查，第二个 SSE frame 容易与结束事件合并，其 TPOT 不作为稳态性能结论。

调度公平性分成两种负载：

8 个 2048-word 长 Prefill 与 8 个 32-word 短请求同时到达时，短请求 TTFT p50 约 30.142 s，虽早于长请求的 44.837 s，但短请求后续 TPOT p50 达 1.764 s。这说明 Chunked Prefill 能让短请求提前产出首 token，但长 Prefill 仍会严重阻塞短请求 Decode，是当前最明确的调度弱点。
8 个 32-word/128-output 请求进入 Decode 后第 3 秒，再注入 8 个 32-word/8-output 请求。短请求空载基线 TTFT/TPOT 为 1.327 s/21.508 ms；动态注入后两轮中位值约为 1.585 s/39.558 ms，吞吐从 43.180 降至 34.403 token/s。但新请求在提交后约 1.862 s 完成，明显早于约 12.2 s 才完成的后台请求，证明 Continuous Batching 可在已有 Decode 中接入请求且没有饥饿；代价是 TTFT +19.4%、TPOT +83.9%、短请求吞吐 -20.3%。

16 路、每路 1024-word/8-output 的 Prefill chunk sweep，每档两轮：

Prefill chunk	TTFT p50	TPOT p50	Prompt words/s	单轮 TPOT p95 最大值
128	38.076 s	45.409 ms	426.610	46.366 ms
256	38.017 s	36.050 ms	425.525	46.560 ms
512	38.173 s	33.211 ms	424.185	659.670 ms

三档 Prompt 吞吐差异不足 0.6%，扩大 chunk 没有带来吞吐收益；512-token chunk 产生 659.670 ms 的 TPOT 尾部离群，调度不均衡最明显。128 的请求完成分布最均匀， 256 的中位 TPOT 更好，因此默认仍保留 256；若 workload 更强调所有请求公平完成， 128 是可选配置，512 不推荐用于当前 16 路负载。

实验驱动为 tools/run_tp4_performance_matrix.py。正式 artifacts：

build-perf/qwen3-14b-tp4/performance-matrix-16x4096.json，SHA256 6af4128bfaea36eb0f10200375982acdf51e514722e60625039dc434468dffe2；
build-perf/qwen3-14b-tp4/performance-matrix-output-arrival.json，SHA256 7442d2945ca98f97dc0902e9ab407e173a805e38e672021c28f987022f6272a9；
build-perf/qwen3-14b-tp4/chunk-probe-128.json，SHA256 7d2e08800ae7492d3699f11312e83a749ea2315227308aa7a7e4b216556bc6ca；
build-perf/qwen3-14b-tp4/chunk-probe-256.json，SHA256 37a4f175d80c8fe239fe93554a30706fbfd3546b1c4ee96cebf6e99a46aaf3e3；
build-perf/qwen3-14b-tp4/chunk-probe-512.json，SHA256 550c6e32577bf9fa21c928eeb8f6a296c323b9fc4fe5dcaabd9cdf5aa52b41e2。

Qwen3-14B TP4 Large-M Prefill 收尾

针对 16×4096 压力实验暴露出的 Prefill 吞吐问题，新增 Rank-local Large-M GEMM benchmark，覆盖 QKV [M,5120] x [1792,5120]、O [M,1280] x [5120,1280]、Gate/Up [M,5120] x [8704,5120] 和 Down [M,4352] x [5120,4352]，并扫描 M=256/512/1024/2048/4096。每个候选执行 5 次 warmup 和 10 次测量，检查采样输出误差与 steady-state allocation。

初始 DTK 4.4 环境只有 hipBLAS/rocBLAS，没有预装 hipBLASLt 或 Composable Kernel；该版本 rocBLAS 的 solution_index 仍为 reserved，不能进行算法枚举。gfx906 上的 v_mfma_f32_32x32x4bf16 汇编探针明确失败，而相同探针以 gfx908 为目标可以汇编和反汇编，因此本轮不能声称启用了 BF16 MFMA。可用候选只有 hipBLAS atomics 开关与固定权重布局。Atomics 没有稳定收益；把权重从 [N,K] 预转置为 [K,N] 后，所有 shape 均保持数值正确，QKV 提升 1.42-1.57x、O 提升 1.15-1.18x、Gate/Up 提升 1.44-1.56x、Down 提升 1.30-1.45x。

生产路径仅为 Qwen3-14B TP4 的 Gate/Up 保存预打包副本，并仅在 Prefill M>=256 时调用 no-transpose hipBLAS；Decode 继续使用原 [N,K] 权重和 small-M kernel。Materialization 使用 32×32 分块 CPU 转置生成每层 [5120,8704] 副本。该选择每卡额外占用约 3.32 GiB，是用容量换 Prefill 吞吐，不能用于其他模型或 TP 规模。Context A/B 结果如下；32/128 words 没有触发 M>=256 路由，变化属于运行波动。

Prompt words	TTFT 变化	Prompt words/s 变化
32	+0.55%	-0.49%
128	+0.17%	-0.16%
512	-10.53%	+11.24%
1024	-10.68%	+11.65%
2048	-9.55%	+10.39%
4080	-7.55%	+8.09%

16 个并发请求、每请求 4088 prompt tokens 的满容量实验仍为 16/16 成功。相对未预打包的 max_batched_tokens=4096 基线，总墙钟由 211.306 s 降至 191.135 s （-9.55%），实际 prompt 吞吐由 309.541 提升至约 342.211 token/s （+10.55%），请求延迟 min/median/max 由 194.455/211.294/211.301 s 降至 175.442/191.121/191.132 s。空载显存从约 11.7 GiB/卡 增至约 15.1 GiB/卡，只余约 0.9 GiB/卡，因此该配置已经接近当前 16 GiB Device 的容量边界。使用 32×32 分块转置后的最终二进制做了一次完整冷启动，约 137 s 后通过 /health，并成功完成 Hello -> ": I have a" 的四 token greedy HTTP 推理；该单次启动时间只用于确认 materialization 没有异常，不作为稳定性能分位数。

第一代 10Q/2KV/head_dim=128 tiled Paged Flash Prefill Attention 只让同一 KV Head 的 5 个 Query Heads 共享 KV16 LDS tile，每个 Query token 仍由一个 wave 串行扫描完整 Context，因此加权几何平均只有 1.03056x，没有进入生产路径。

第二代 kernel 同时沿 Query 与 KV 两个维度分块。一个 320-thread block 仍由 5 个 waves 覆盖 5 个 GQA heads，但每个 wave 被划分成 8 个独立 8-lane subgroup；每个 subgroup 并行负责一个 Query token，每 lane 处理 16 个 head-dim 元素。KV16 以 BF16 pair 形式一次读入 8260 B LDS，并在 5 heads × 8 queries 之间复用。每个 Query 保持独立的 causal mask 和 online-softmax 状态，不建立 score/packed workspace。序列内以 QueryTile 8 对齐，尾部不足 8 token 时只启用有效 subgroup，因此不会跨 ragged sequence 共享 KV。

有限候选扫描中，Q2/KV16、Q4/KV32、Q8/KV16 的加权几何平均分别约为 1.070x/1.448x/1.883x；串行保存四个 Query 状态的早期 Q4 实现因丢失 Query 并行度只有 0.453x。最终 Q8/KV16 code object 使用 73 VGPR / 70 SGPR / 8260 B LDS，无 VGPR、SGPR 或 scratch spill。

Q8/KV16 对六个场景执行五轮正式 3 warmup / 10 iterations 测试。五轮加权几何平均范围为 1.8693-1.8729x，全部超过 1.50x 生产门槛；single-4096 为 1.8116-1.8315x，batch16-late 为 1.8862-1.8943x。最短的 single-256 只有 1.1282-1.1671x，因此生产路由只对白名单 10Q/2KV/head_dim=128/block16 且平均每条 active sequence 本轮 Prefill tokens >=256 时启用，其余场景保留 scalar fused kernel。

subgroup reduction 改变了 FP32 加法顺序，输出不再逐位相同；六场景最大绝对误差为 0.000244140625，最大相对误差为 0.0078125，均通过既有 BF16 数值门槛。额外的 3+509 ragged sequence 边界测试验证了 QTile 尾部、跨序列隔离、K/V append 和生产路由选择。

本阶段 artifacts：

build-perf/qwen3-14b-tp4/large-m-gemm-autotune.json，SHA256 0293255268cefa6290ac51b307f8c645e5ea8f9a952729d11ac0ef40476eb6f6；
build-perf/qwen3-14b-tp4/prefill-attention-tiled-kv16-final.json，SHA256 c96099e69984fa3f615460f882a8833b972b81399ee4174ae8e1415049d1c2fa；
build-perf/qwen3-14b-tp4/prefill-attention-subgroup-q8-kv16-round-1.json， SHA256 ed8d1b3ecfae5063c824421e25bda2f61e7a0c9c2cc41bec165aa5376a91bc76；
build-perf/qwen3-14b-tp4/prefill-attention-subgroup-q8-kv16-round-2.json， SHA256 501b4e7d920ccb4f889b9afab1fb8e2c9855d0d19d3fde78fbfae5fc2730039e；
build-perf/qwen3-14b-tp4/prefill-attention-subgroup-q8-kv16-round-3.json， SHA256 7f39f677e9bc0f258a32ec9b10f4e377406bef46ff154493004d0b70d1040d8b；
build-perf/qwen3-14b-tp4/prefill-attention-subgroup-q8-kv16-round-4.json， SHA256 2535896d0b9e1cd9316558757c2361265b40621da40e07d8702b0706cd0c3994；
build-perf/qwen3-14b-tp4/prefill-attention-subgroup-q8-kv16-round-5.json， SHA256 c4c228e98c11b4e9992420fecfdd22b6c5e7174834e2a1706ee2b707cb42542b；
build-perf/qwen3-14b-tp4/prepacked-gate-up-context-matrix.json，SHA256 ca3bea502706aa7cbf5bd1962ef71ca9197b8975054ffe51c59f201ed585eb02；
build-perf/qwen3-14b-tp4/prepacked-gate-up-capacity-16x4096.json，SHA256 626624ca76d29cd6df28e0e4a0e51035f52743b1aba44557aaadc3d9618ec928。

最终验证包括完整 HIP build、通用 CTest 141/141、HIP Backend 12/12、HIP Operators 24/24、Dynamic Engine HIP 10/10、Tiny HIP 6/6、TP collective 9/9、TP4 RCCL 2/2，以及本阶段 GEMM/Attention tooling 2/2，全部通过。

Tiled Prefill Attention 端到端 A/B

为避免把 Attention microbenchmark 的 1.87x 直接外推为整模型收益，同一正式服务二进制增加了 --tiled-prefill-attention 0|1 实验开关，默认值为 1。A/B 两侧只切换 Qwen3-14B TP4 Prefill Attention 的 scalar/tiled 路由；模型、权重预打包、TP4 RCCL、 16 active x 4096 context KV 容量、max_batched_tokens=2048、256-token chunk、采样和 Decode 策略保持一致。每个服务完成权重加载后先执行固定短请求预热，再运行 Context 六档各两轮和 16 x 4080 words 容量压力一轮。

单请求 Context 结果如下。吞吐使用既有矩阵的 prompt word-count proxy；32/128 words 没有触发 tiled 路由，结果持平。随着 Context 增长，Attention 占比上升，端到端收益逐步接近但仍低于纯 Attention operator 加速。

Prompt words	Scalar TTFT	Tiled TTFT	TTFT 变化	Scalar words/s	Tiled words/s	吞吐变化
32	434.650 ms	434.658 ms	+0.00%	66.413	66.402	-0.02%
128	495.927 ms	495.755 ms	-0.03%	234.744	234.837	+0.04%
512	1266.110 ms	1249.042 ms	-1.35%	386.529	391.617	+1.32%
1024	2540.943 ms	2473.119 ms	-2.67%	391.928	402.654	+2.74%
2048	5703.403 ms	5378.367 ms	-5.70%	353.277	374.266	+5.94%
4080	14108.148 ms	13002.091 ms	-7.84%	286.356	310.465	+8.42%

容量压力中，16 个请求各自实际包含 4088 prompt tokens，总计 65,408 tokens。总墙钟从 193.355 s 降至 160.202 s（-17.15%），TTFT p50 从 141.666 s 降至 116.816 s（-17.54%），TTFT min/max 从 89.328/193.341 s 降至 72.784/160.188 s。Prompt 吞吐从 337.617 提升到 407.486 words/s （+20.69%）；按实际 tokenizer token 数计算，对应 338.279 -> 408.285 token/s。 16/16 请求均完成，没有 OOM、KV exhaustion 或 RCCL abort，服务退出后四卡显存均回到约 2 MiB。

共享 Native Profiler 对两侧记录了完全相同的 29 个请求、325,108 rank-summed executed tokens 和 245,360 个 operator events，均无 dropped/error event。四 Rank 累计 Prefill graph 时间由 961.877 s 降至 817.170 s（-15.04%），Decode 累计时间仅变化 -0.30%。容量主路径的 2048 tokens x 8 sequences Prefill step 在每个 Rank 上共记录 112 次，平均整图时间由 5.750 s 降至 4.567 s（-20.58%）；混合 Context 的 256 x 1 step 由 715.971 ms 降至 664.709 ms（-7.16%）。未进入白名单的 248 x 1、2040 x 8 和 2040 x 9 基本持平，证明收益来自 tiled Prefill 路由而非 Decode 或服务测量波动。

本次 A/B artifacts：

build-perf/qwen3-14b-tp4/prefill-tiled-ab-scalar-context.json，SHA256 57893830d343fde5a572e31807341c974695aa16210e8f739662e1de0ebc01f6；
build-perf/qwen3-14b-tp4/prefill-tiled-ab-tiled-context.json，SHA256 269baabbd7171af2d3de8c93e03d52a60ca9450acf48e0b3fc6b9ba3210a7d5a；
build-perf/qwen3-14b-tp4/prefill-tiled-ab-scalar-capacity.json，SHA256 4fe77a39dcf8832816734a8ed51ae54118db6af81a2152c218ef6223e6df5153；
build-perf/qwen3-14b-tp4/prefill-tiled-ab-tiled-capacity.json，SHA256 7fccb5f8e1a68d044f65ca62228d6e43a050bea4b9416639e98a273b28cb72b6；
build-perf/qwen3-14b-tp4/profile-prefill-scalar-e2e/{metrics,trace}.json， SHA256 84e6daa228ee138ff2f4b5d9229c8b8fced73ed8f72a6ce0c4e051cfafd36e04 / db77798ee53a6d0b5a314bfb94043bd13a0be129a387f7977458cab51e90e365；
build-perf/qwen3-14b-tp4/profile-prefill-tiled-e2e/{metrics,trace}.json， SHA256 2943ff1b477a79becca319a3374534219932a5992430381cbd0315da58c2663d / a558901cfc82ede2b1545048149c5ac8bffe3a9581c9a5c755d63f915dc08ad2。

hipBLASLt / Composable Kernel 实验接入

两套外部 GEMM 项目以 ROCm 5.7.1 对应 revision 和 archive/tree SHA256 固定在 tools/external_gemm_dependencies.json，由 tools/sync_external_gemm_dependencies.py 下载到用户缓存，不加入正常 runtime vendor tree。这样默认构建不增加源码体积，也不会让实验依赖影响生产服务。

python3 tools/sync_external_gemm_dependencies.py --component all

cmake -S . -B build-external-gemm \
  -DCMAKE_C_COMPILER=/opt/dtk/llvm/bin/clang \
  -DCMAKE_CXX_COMPILER=/opt/dtk/llvm/bin/clang++ \
  -DCMAKE_HIP_COMPILER=/opt/dtk/llvm/bin/clang++ \
  -DMETAINFER_REF_ENABLE_HIP=ON \
  -DMETAINFER_REF_ENABLE_EXTERNAL_GEMM_PROBES=ON \
  -DMETAINFER_REF_CK_ROOT="$HOME/.cache/metainfer/external-gemm/src/composable_kernel"

cmake --build build-external-gemm \
  --target metainfer_ref_ck_bf16_gemm_probe
build-external-gemm/bin/metainfer_ref_ck_bf16_gemm_probe 256 1792 5120

ROCm 5.7.1 的 CK 官方 BF16 GEMM 实例全部基于 XDL，运行时只接受 gfx908/gfx90a/gfx940，因此 gfx906 会返回 unsupported。实验 probe 为通用 DeviceGemmDl 补了 CK 5.7 未提供的 bhalf2 × bhalf2 -> FP32 普通向量乘加；该路径可以运行，但不使用 MFMA。项目 target 在纯净 CK 源码上编译通过，QKV [M=256,K=5120,N=1792] 实测为 2.304 ms / 2.039 TFLOPS。

与现有预打包 hipBLAS 的代表点比较，CK DL 在 M=256 四种投影慢 1.35-1.72x，在 M=4096 慢 1.10-1.50x。其中 M=4096 Gate/Up 为 99.419 ms，现有 hipBLAS 为 66.107 ms。因此 CK 只保留为能力和性能探针，不进入 LinearBf16 生产路由。

hipBLASLt 5.7.1 的 API/epilogue host 层可以在当前 toolchain 编译，但随包 180 个 Tensile logic 只覆盖 gfx90a 与 gfx940/941/942，没有 gfx906 算法。关闭 Tensile 构建得到的 host library 仍依赖算法后端符号，不能作为 GEMM fallback；完整安装则必须同时提供匹配架构的 Tensile library。metainfer_ref_hipblaslt_bf16_gemm_probe 已接入 hipblasLtMatmulAlgoGetHeuristic，供存在完整 hipBLASLt 安装或更换受支持 GPU 后直接查询 algorithm count，但当前服务器不启用该 target，也不改变生产路由。

Qwen3-14B TP4 B8/B16 Decode 收尾

Decode-aware Prefill 公平调度之后，Rank-specific Timeline 显示 B8/B16 的主要设备时间仍为 GEMM 和 RCCL。新增 BF16x4 medium-M Linear 精确白名单后，B8/B16 四类 projection microbenchmark 提升约 1.34-2.24x；生产路径只覆盖 Qwen3-14B TP4 的精确 B8/B16 shape。Persistent HTTP 并发基线为 B8 59.166 token/s / 52.058 ms TPOT，B16 71.082 token/s / 84.565 ms TPOT。

随后对 RCCL 做了逐层裁决。[8,5120] BF16 在孤立 benchmark 中拆成两个 40 KiB collective 可由 107.974 us 降至 71.857 us，但真实模型每步物理 AllReduce 从 80 次增加到 160 次，四 Rank collective family 由 16.542-21.897 ms 增至 34.268-48.099 ms，因此完整回退。全局强制 LL128 可把 B8/B16 分别降至 50.618/72.809 us，但 256-row Prefill 从 387.182 us 恶化到 941.192 us； Tree、Tree+LL128 和自动 LL128 也未通过门槛。四卡 peer scratch 加单 kernel 的 hybrid P2P 候选数值正确，但跨 NUMA 3-hop PCIe 上 B8/B16 为 122.202/171.573 us，分别比单次 RCCL 慢 13.18%/43.93%。当前 DTK RCCL 没有 tuner plugin 或 per-communicator 协议 API，所以生产路径继续使用单次 RCCL，不保留上述候选代码。

Elementwise 检查发现 BF16x2 paired AddRmsNorm 的 kernel 可按行独立执行，但旧路由只允许 B1/B4。设备测试和 benchmark 扩展到 B8/B16 后，五轮 20 warmup / 200 iterations 均满足 updated residual 逐位一致、normalized 零误差和 allocation delta 为 0。五轮加速范围/中位数如下：

Shape	加速范围	中位加速
B1 x 5120	1.7374-1.8157x	1.8019x
B4 x 5120	1.6718-1.7103x	1.6954x
B8 x 5120	1.6835-1.6952x	1.6911x
B16 x 5120	1.6714-1.6805x	1.6750x

生产路由因此从 rows<=4 放宽到 rows<=16，仍严格要求 hidden=5120、BF16、连续且 BF16-pair 对齐。新 Timeline 中 B8/B16 每步均为 80 次 AddRmsNormPairsKernel，elementwise family 从 4.192-4.256 ms 降至 2.775-2.838 ms，减少 32.4%-34.5%，GEMM 和 Attention 结构不变。

同一 16 active x 4096 context persistent 服务三轮并发矩阵结果：

并发	TTFT 基线 -> 新值	TPOT 基线 -> 新值	吞吐基线 -> 新值
B1	495.480 -> 495.603 ms	20.494 -> 20.494 ms	28.263 -> 28.278 token/s
B4	1269.539 -> 1270.802 ms	23.729 -> 23.720 ms	63.779 -> 63.742 token/s
B8	2708.646 -> 2699.752 ms	52.058 -> 50.827 ms	59.166 -> 59.691 token/s
B16	4572.896 -> 4573.830 ms	84.565 -> 83.534 ms	71.082 -> 71.395 token/s

B8 TPOT/吞吐改善 2.36%/0.89%，B16 改善 1.22%/0.44%；B1/B4 基本持平。该优化正式保留，但 B8 尚未达到 65 token/s / 45 ms，B16 TPOT 也尚未达到 80 ms，后续收益仍需来自 GEMM 或能真正按消息选择协议的通信后端，而不是继续堆叠 elementwise launch fusion。

本阶段 artifacts：

build-perf/qwen3-14b-tp4/b8-b16-elementwise-round-{1..5}.json；
build-perf/qwen3-14b-tp4/kernel-timeline-b8-b16-paired-elementwise.json，SHA256 424390a993a630eda57ccb5eabd3feda93499f7bb657487d78931112e1f9d3ab；
build-perf/qwen3-14b-tp4/decode-b8-b16-paired-elementwise-concurrency.json， SHA256 6d7058e1895cf9e127aadfbd412a0727c573ee2e902fa210f8686def4f2dcd86。

最终验证包括完整 Release build、通用 CTest 142/142、HIP Backend 12/12、 HIP Operators 24/24、Dynamic Engine HIP 10/10、Tiny HIP 6/6、TP2 collective 9/9、TP4 RCCL 3/3 和受影响 tooling 4/4，全部通过；服务退出后四卡均回到约 2 MiB。

B8/B16 LDS Input-Reuse GEMM

上一轮 BF16x4 medium-M kernel 中，同一个 256-thread block 的四个 wave 分别计算四个 output row，但每个 wave 都会从全局显存重复读取相同的 B8/B16 input。新的 PackedMediumMQuadLdsBf16LinearKernel 把 K 维切成 BF16 quad tile；每个 block 先把一个 token group 的 input tile 合作搬入 LDS，再由四个 wave 共享。kernel 仍使用 FP32 累加并保持原累加顺序，不引入运行时 workspace 或额外 allocation。

候选同时扫描 rows_per_group=8/16 和 input_tile_quads=64/128/256。五轮正式测试使用 20 warmup / 200 iterations，所有候选均数值正确且 allocation delta 为 0。相对当时生产路径的逐轮最差/中位加速如下：

Shape	生产选择	最差加速	中位加速
B8 QKV	group8, tile128	1.9720x	1.9785x
B8 Gate/Up	group8, tile128	1.5911x	1.5919x
B8 Down	group8, tile128	1.3717x	1.3755x
B16 QKV	group8, tile128	1.6149x	1.6182x
B16 Gate/Up	group8, tile128	1.4364x	1.4379x
B16 Down	group8, tile128	1.2907x	1.2917x

O projection 没有进入 LDS 白名单。B8 的最好 LDS 候选只有约 1.02x，低于 3% 生产门槛；B16 的最好 LDS 候选仍慢于已有 group16 kernel。最终自动路由只覆盖 Qwen3-14B TP4 的精确 B8/B16 QKV、Gate/Up、Down shape，并固定为 group8 + tile128；O projection、其他 batch 和其他模型继续使用原 measured policy。

Rank-specific Timeline 中，四个 Rank 的 B8 GEMM family 总时间下降 20.9%-21.5%，B16 下降 20.8%-21.0%。B8 range wall time从 148.204 ms 降至 137.090 ms，下降 7.50%。B16 该次 profiler 采样的 RCCL 时间异常升高，因此不以单次 Timeline wall time裁决，而以无 profiler 的 persistent HTTP 三轮中位数作为端到端结果：

并发	TTFT 旧值 -> 新值	TPOT 旧值 -> 新值	吞吐旧值 -> 新值
B1	495.603 -> 503.460 ms	20.494 -> 20.520 ms	28.278 -> 27.793 token/s
B4	1270.802 -> 1268.918 ms	23.720 -> 23.603 ms	63.742 -> 63.922 token/s
B8	2699.752 -> 2636.629 ms	50.827 -> 41.734 ms	59.691 -> 65.069 token/s
B16	4573.830 -> 4481.684 ms	83.534 -> 71.358 ms	71.395 -> 76.403 token/s

B8 TPOT/吞吐改善 17.89%/9.01%，B16 改善 14.58%/7.01%，达到 B8 >=65 token/s / <=45 ms 和 B16 >=70 token/s / <=80 ms 的阶段目标。 B1 吞吐回退 1.71%，仍在 2% 门槛内；B4 基本持平。

本阶段 artifacts：

build-perf/qwen3-14b-tp4/decode-b8-b16-quad-lds-round-{1..5}.json；
build-perf/qwen3-14b-tp4/kernel-timeline-b8-b16-quad-lds.json，SHA256 c4b60b81790fac54e5bf79cbf24c8e5ee5167092ddc96919d1aced84a95250fe；
build-perf/qwen3-14b-tp4/decode-b8-b16-quad-lds-concurrency.json，SHA256 18a1fa9a47f626356c35c9490492f9e7fd9621c3026ee53ff7c368f732203362。

Qwen3-14B TP4 W8A16 Decode 与 Selective Prefill

当前量化格式为 signed INT8 per-output-channel weight、BF16 activation、F32 scale 和 FP32 accumulation。Embedding 与 Norm 保持 BF16；Q/K/V/O、Gate/Up/Down 和 LM Head 使用对称 [-127,127] 权重量化。INT4 尚未实现。

离线生成和独立校验：

build-perf/bin/metainfer_ref_int8_quantize \
  --model-dir "$HOME/.cache/metainfer/models/Qwen3-14B" \
  --output build-perf/int8-qwen3-14b/model.int8.safetensors

build-perf/bin/metainfer_ref_int8_quantize \
  --model-dir "$HOME/.cache/metainfer/models/Qwen3-14B" \
  --check build-perf/int8-qwen3-14b/model.int8.safetensors

真实 14B artifact 将 29,536,614,400 字节源数据降为 15,555,608,064 字节（52.67%），包含 281 个量化 tensor 和 162 个 BF16 copy tensor；SHA256 为 06856d9451a3aa4426b31dac8dc44c83f4ac065c991c9101ca95155649fb93f6。Artifact 保存在本地 build 目录，不进入 Git。

生产服务使用双权重模式。BF16 权重始终存在，INT8 仅对五轮 microbenchmark 达到 >=1.25x 的精确 shape/batch 路由；未通过门槛的 shape/M 继续走 BF16：

路径	INT8 batch	五轮中位加速
QKV	B1 / B4	1.482x / 1.290x
O	B1	1.262x
Gate/Up	B1	1.840x
Down	B1	1.607x
TP4 F32 LM Head	B1 / B4 / B8 / B16	10.365x / 3.777x / 2.747x / 1.396x

M>16 Prefill 对 RowsPerGroup=8/16 和 input_tile_quads=64/128/256 做了扫描，最终仍选择 group8 + tile128。五轮 M32/M64 的中位加速如下；M128/M256/M512 未达到门槛并明确回退 BF16：

路径	M32	M64
QKV	2.919x	2.994x
O	1.588x	1.376x
Gate/Up	3.095x	1.735x
Down	2.161x	1.830x

五轮 benchmark 使用 20 warmup / 200 iterations：

build-perf/bin/metainfer_ref_int8_decode_linear_benchmark \
  --hip-device 0 --warmups 20 --iterations 200 --rounds 5 \
  --output build-perf/benchmarks/qwen14-tp4-int8-decode-linear.json

build-perf/bin/metainfer_ref_int8_decode_linear_benchmark \
  --prefill --hip-device 0 --warmups 20 --iterations 200 --rounds 5 \
  --output build-perf/benchmarks/qwen14-tp4-int8-prefill-linear.json

Qwen3-14B TP4 服务启动：

build-perf/bin/metainfer_ref_qwen3_server \
  --model-dir "$HOME/.cache/metainfer/models/Qwen3-14B" \
  --int8-weights build-perf/int8-qwen3-14b/model.int8.safetensors \
  --devices 0,1,2,3 --host 127.0.0.1 --port 8000

4x16 GiB gfx906 实机物化后占用约 13.98-14.03 GiB/卡。固定 1-token prompt、 2-token greedy generation 的五次暖服务中位数由 BF16 45.81 ms 降至 INT8 28.97 ms，约 1.58x，且两侧生成 token 完全一致。当前服务 INT8 路由只接受 Qwen3-14B TP4 BatchService。

固定 32-token prompt + 2-token greedy generation 的五次暖服务中位数由 BF16 302.60 ms 降至 INT8 132.84 ms，约 2.28x；64-token prompt 则由 404.27 ms 降至 229.77 ms，约 1.76x。两组 A/B 的 greedy 输出均完全一致。其他模型、单卡和 TP2 的 INT8 KV 路由仍保持未启用；14B TP4 的 M256/M1024/M4096 Prefill 已通过 Hybrid 路径验证。

INT8 Paged KV Cache 与 Attention

Paged KV Cache 支持可选的 signed INT8 K/V pool。K 和 V 分别使用 per-token-per-KV-head FP32 scale，append 时在线对称量化到 [-127,127]；Decode 从 INT8 page 读取并反量化。Prefill 使用 Hybrid 路径：历史位置从 INT8 page 读取，当前 chunk 直接读取尚未离开算子输入的 BF16 K/V，并在 attention 前将当前 K/V 量化写入 Cache，供下一个 chunk 和 Decode 使用。这样不会把当前 chunk 做一次“BF16 -> INT8 -> BF16”的无效往返，也不创建完整 BF16 cache 或 score workspace。逻辑 block table、reservation、rollback 和 generation 语义与 BF16 cache 完全相同。

每个 token/head 的 BF16 K+V 需要 2 * 128 * 2 = 512 bytes；INT8 K+V 加两个 FP32 scale 需要 2 * 128 + 2 * 4 = 264 bytes，因此固定布局容量比为 0.515625，即 KV 数据区减少 48.4375%。Cache 的 K/V 与 scale 仍由一次对齐分配提供，不增加 steady-state allocation。

服务默认保持 BF16。实验性启用方式：

build-perf/bin/metainfer_ref_qwen3_server \
  --model-dir "$HOME/.cache/metainfer/models/Qwen3-14B" \
  --int8-weights build-perf/int8-qwen3-14b/model.int8.safetensors \
  --devices 0,1,2,3 --int8-kv-cache 1 \
  --host 127.0.0.1 --port 8000

独立 Attention benchmark 固定为 Qwen3-14B TP4 本地 10Q/2KV/128 shape，Decode 覆盖 B1/B4/B8/B16、4096 context，Prefill 覆盖 256/1024/4096 tokens：

build-perf/bin/metainfer_ref_qwen14_int8_kv_attention_benchmark \
  --hip-device 0 --warmups 3 --iterations 10 \
  --output build-perf/benchmarks/qwen14-int8-kv-hybrid-prefill.json

Hybrid Prefill 和 INT8 Decode 的五轮正式测试中位结果如下。测试固定为 Qwen3-14B TP4 本地 10Q/2KV/128 shape，Prefill case 是一次性 Prefill（past_length=0），因此它主要验证当前 chunk 不再经过量化往返；在更小 chunk 的长上下文场景中，历史部分仍会承担 INT8 反量化开销。speedup < 1 表示 INT8 KV 更慢：

Case	BF16	Hybrid INT8 KV	speedup
Decode B1 / 4096	1.131 ms	1.201 ms	0.942x
Decode B4 / 4096	1.207 ms	1.481 ms	0.818x
Decode B8 / 4096	1.559 ms	1.817 ms	0.858x
Decode B16 / 4096	2.148 ms	2.411 ms	0.891x
Prefill 256	0.778 ms	0.777 ms	0.977x
Prefill 1024	5.748 ms	5.996 ms	0.949x
Prefill 4096	77.849 ms	79.884 ms	0.973x

5/5 轮数值与 allocation 检查通过。Decode 最大绝对误差为 7.63e-6，Prefill 最大绝对误差为 4.88e-4。1-token prompt + 2-token greedy 的 TP4 暖服务 A/B 输出完全一致；中位端到端延迟约从 BF16 KV 的 28.64 ms 变为 INT8 KV 的 29.08 ms。

统一 BF16 tile 后，Decode 几何平均为 0.876x，Hybrid Prefill 为 0.964x。相对于旧的“先量化再从 INT8 读取当前 chunk”实现，Prefill 几何平均从 0.740x 提升到 0.964x，4096-token case 从约 0.725x 恢复到 0.973x；五轮 correctness 与 allocation 检查全部通过。Decode 仍受 INT8 unpack/FP32 dequant 成本影响，未通过性能 gate，因此 INT8 KV 保持显式实验开关，不作为默认生产路由；它的主要收益仍是固定显存预算下约 48.4% 的 KV 容量节省。

Qwen3-32B 双节点 TP4 x PP2 短期版本

短期分布式拓扑固定为两个节点、每节点四张 GPU。Node 0 的四个 Rank 组成第一个 TP4 Group，执行 Embedding 和 Layer 0-31；Node 1 的四个 Rank 组成第二个 TP4 Group，执行 Layer 32-63、FinalNorm、Vocab-parallel LM Head 和 Sampling。每个 Stage 的 Row-parallel AllReduce 已使四个 TP Rank 持有完整 BF16 hidden state，因此跨节点数据面只由 root pair 0 <-> 4 传输一份 activation。Rank 4 接收后通过 Node 1 TP Broadcast 分发，token/valid 也只由 Rank 4 返回 Rank 0，再通过 Node 0 TP Broadcast 分发。Node 0 是唯一的 HTTP/OpenAI Endpoint，Node 1 只运行远端 Stage Worker。

节点间数据面使用 RCCL，控制面和 communicator bootstrap 使用 TCP。实机验证环境为 worker5/11.8.2.48 与 worker4/11.8.2.47，RCCL 日志已确认选择 NET/IB 和 mlx5_0:1，即走 ibp1s0 对应的 200 Gb/s InfiniBand。两节点的实际 RCCL DMA-BUF 测试均打印 GPU Direct RDMA Enabled for HCA 0 'mlx5_0'，并通过四卡 TP AllReduce 与跨节点 PP activation 的数值校验。这里使用的是 DTK/HSA 的 hsa_amd_portable_export_dmabuf + ibv_reg_dmabuf_mr 路径，不是旧式 PeerDirect 注册路径。

因此，旧的 ibv_reg_mr(HIP pointer) 返回 Bad address 不能作为 GDR 不可用的依据：在本机 hydcu 驱动上它只代表 legacy PeerDirect 路径不可用。项目的 metainfer_ref_gpudirect_probe 已优先验证 DMA-BUF 注册；RCCL 日志中的 GPU Direct RDMA Enabled 是服务实际是否启用 GDR 的最终判据。不要为旧探针失败安装 nvidia-peermem、重装 OFED，或盲目修改 IOMMU/ACS。

构建和模型前提

cmake --build build-perf --target metainfer_ref_qwen3_distributed_node -j4

build-perf/bin/metainfer_ref_gpudirect_probe --device 0 --hca mlx5_0

该探针应输出 GPUDirect RDMA registration available。当前仅 mlx5_0:1 的 ibp1s0 链路为 Up；mlx5_1..3 对应端口尚未连通，因此服务应继续固定使用 mlx5_0:1，不能在没有链路和拓扑验证前将四张 GPU 强行分配到其他 HCA。

两个节点必须安装兼容的 /opt/dtk Runtime，运行用户必须属于 render 和 video Group，并且四张 GPU 均可见。短期版虽然每个节点只把本 Stage 的 32 层和本地 TP Shard 物化到 GPU，但 WeightArchive::Open 仍校验完整 Safetensors Archive。因此两个节点都必须通过本地磁盘、NFS 或其他共享挂载看到同一个完整 Qwen3-32B 模型目录。BF16 完整权重约 64 GiB，不能把只有一半 shard 的目录直接交给当前 Worker。

启动前在两个节点检查：

id
test -r "$MODEL_DIR/config.json"
test -r "$MODEL_DIR/model.safetensors.index.json"
hy-smi --showmeminfo vram | head -n 20

启动顺序

先在 Node 0 worker5 启动协调器。--coordinator 同时是 bootstrap/control 的监听地址，必须使用 Node 1 可达的地址；不要写 127.0.0.1。

export MODEL_DIR=/shared/models/Qwen3-32B
export LD_LIBRARY_PATH=/opt/dtk/lib:/opt/dtk/hip/lib:${LD_LIBRARY_PATH:-}
export NCCL_SOCKET_IFNAME=ibp1s0
export NCCL_IB_HCA=mlx5_0:1
export NCCL_IB_DISABLE=0
export NCCL_DMABUF_ENABLE=1
export NCCL_NET_GDR_LEVEL=5

build-perf/bin/metainfer_ref_qwen3_distributed_node \
  --model-dir "$MODEL_DIR" \
  --node-rank 0 --coordinator 11.8.2.48 \
  --bootstrap-port 29500 --control-port 29510 \
  --devices 0,1,2,3 --host 0.0.0.0 --port 8000 \
  --max-context-tokens 4096 --max-output-tokens 256 \
  --max-active-sequences 8 --max-batched-tokens 2048 \
  --prefill-chunk-tokens 256 --max-mixed-prefill-tokens 16 \
  --admission-coalescing-us 2000

需要在启动前强制验证本机每张 GPU 都能注册 DMA-BUF GDR 时，在两个节点追加：

--require-gpudirect 1 --gpudirect-hca mlx5_0

Worker 会在加载模型权重前直接测试 GPU buffer 的 DMA-BUF RDMA registration。只有探针成功且 RCCL 以 NCCL_DEBUG=INFO NCCL_DEBUG_SUBSYS=NET 启动时显示 GPU Direct RDMA Enabled，才应把该配置作为 GDR 服务基线。

随后在 Node 1 worker4 使用相同模型路径和容量参数启动后半模型：

export MODEL_DIR=/shared/models/Qwen3-32B
export LD_LIBRARY_PATH=/opt/dtk/lib:/opt/dtk/hip/lib:${LD_LIBRARY_PATH:-}
export NCCL_SOCKET_IFNAME=ibp1s0
export NCCL_IB_HCA=mlx5_0:1
export NCCL_IB_DISABLE=0
export NCCL_DMABUF_ENABLE=1
export NCCL_NET_GDR_LEVEL=5

$HOME/metainfer-shortterm/bin/metainfer_ref_qwen3_distributed_node \
  --model-dir "$MODEL_DIR" \
  --node-rank 1 --coordinator 11.8.2.48 \
  --bootstrap-port 29500 --control-port 29510 \
  --devices 0,1,2,3 --host 127.0.0.1 --port 8000 \
  --max-context-tokens 4096 --max-output-tokens 256 \
  --max-active-sequences 8 --max-batched-tokens 2048 \
  --prefill-chunk-tokens 256 --max-mixed-prefill-tokens 16 \
  --admission-coalescing-us 2000

两边必须使用完全相同的 Context、Batch、Chunk、KV dtype 和模型配置。对于 Qwen3-32B TP4×PP2，--max-mixed-prefill-tokens 16 是当前验证过的默认值：当有 Decode 请求时，它是所有 Prefill 行共享的严格总 token 上限，不会随 Decode 请求数放大。TP4×PP2 Worker 默认还使用 --admission-coalescing-us 2000：仅当引擎完全空闲且已有请求到达时，最多等待 2 ms 合并同一批并发 HTTP 请求的首个 Prefill tick；已有 Decode 不会因此延迟。两节点必须使用相同值，因为它属于 bootstrap contract；可显式传 0 关闭。Node 0 在 bootstrap、RCCL 初始化、权重物化和远端控制握手全部完成后才开放 8000 端口。默认启动等待上限为 300 秒，可用 --startup-timeout-ms 调整。

HTTP/OpenAI 验证

curl -sS http://127.0.0.1:8000/health

curl -sS http://127.0.0.1:8000/v1/chat/completions \
  -H 'Content-Type: application/json' \
  -d '{"model":"Qwen3-32B","messages":[{"role":"user","content":"Hello"}],"max_tokens":4,"temperature":0}'

关闭时只向 Node 0 发送一次 SIGINT 或 SIGTERM。Node 0 先停止 HTTP Admission，再 Drain/Cancel Runtime，并通过控制连接关闭 Node 1；Node 1 不对外开放 HTTP 服务。

当前验收边界

固定 Rank 映射、Stage-local 权重/KV、分阶段 Tiny Qwen3 数值对照、RCCL activation 与结果回传、TCP 生命周期镜像、八进程跨节点通信和“每节点一进程管理四卡”的 Node Smoke 均已通过。2026-07-27 已使用官方 revision 9216db5781bf21249d130ec9da846c4624c16137 完成真实 Qwen3-32B BF16 端到端验收。两个节点分别在 /dev/shm/Qwen3-32B 保存完整 17-shard Archive，共 65,524,328,560 bytes；每个 shard 均已按官方 LFS SHA256 校验。该路径位于 tmpfs，节点重启后模型会消失，不应视为持久模型存储。

真实服务验证结果如下：

/health 和 /v1/models 均返回 Qwen3-32B 正常状态；每卡显存约 10.1 GiB。
连续三次相同 4-token greedy Chat 请求均返回 <think>\nOkay,，HTTP 200；耗时分别为 1.490 s、1.155 s、1.157 s。
B2 同时执行 Chat 与 Completion 时，分别正常返回 <think>\nOkay, the user 和 Paris. True or False，没有跨请求 token 污染。
SSE 能逐 token 输出并以 [DONE] 结束。客户端在 0.8 s 中断长请求后，Cancel 能镜像到远端 Stage，随后请求仍以 HTTP 200 正常返回。
只向 Node 0 发送一次 SIGINT 时，Node 0 会关闭 HTTP admission，并通过控制连接让 Node 1 正常退出。

固定 9-token Chat prompt、16-token greedy 输出的五轮 SSE 基线为：

指标	五轮均值
TTFT	`878.646 ms`
TPOT	`91.261 ms`
单请求端到端生成吞吐	`7.119 token/s`
请求总耗时	`2.248 s`

首次真实服务测试暴露的 compute stream 与 RCCL pipeline stream 竞争已经修复：发送端用 HIP event 让通信流等待 producer，接收端用 event 让 compute stream 等待 activation，不再执行 Host stream synchronize。Pipeline activation 使用两个独立槽；全中间 Prefill chunk 会跳过 LM Head/Sampling 和结果回传，控制面最多保留一拍远端 Tick，在 Stage 1 计算 chunk N 时允许 Stage 0 计算 chunk N+1。最终 Prefill chunk 和 Decode 仍当拍确认。

Root-only、event dependency 和一层深度 Prefill overlap 已使用真实 Qwen3-32B 验证： 562-token prompt 在 128-token chunk 下完成多轮交错 Prefill并返回 HTTP 200；随后短请求和 316-token/9-token B2 均通过，未出现控制响应、KV 或请求间 token 错位。

因此当前状态可以称为“真实 Qwen3-32B TP4 x PP2 双节点短期服务已端到端验收”。仍未完成的长期项包括持久/共享模型存储、将一层 Prefill overlap 扩展到更一般的异步调度，以及与成熟框架在相同 workload 下的横向性能基准。当前 mlx5_0:1 已验证 DMA-BUF GPUDirect RDMA；其余 HCA 端口连通后仍需单独做拓扑与性能验证。

Qwen3-32B Gate/Up 紧凑预打包的容量与性能边界

--prepack-large-m-gate-up 会为每个本 Stage 的 Gate/Up 权重保留一个 [K,N] 副本，供 medium/large-M Prefill 的 no-transpose GEMM 使用。再叠加 --compact-prepacked-gate-up 时，原始 [N,K] GPU 权重会在预打包完成后释放。每个 TP4 Rank 的 Gate/Up shard 为 [12800,5120] BF16；32 个本 Stage 层合计可释放约 3.91 GiB。这是一条显存容量策略，不是默认性能策略。

紧凑格式会失去原 [N,K] B1 Decode 专用 GEMV 的连续读取。保留的紧凑 Decode 路由为：

默认 wave-o8：K=16,N=8 的 wave register-transpose，使用 DS cross-lane 交换，不使用 LDS barrier；
显式 lds16：旧的 K=16,N=32 LDS 路径，仅用于同一二进制 A/B。

在相同二进制、562-token prompt、128-token output 的五轮 B1 对照中，wave-o8 的 TTFT/TPOT/端到端 token/s 为 1875.64 ms / 130.52 ms / 6.94，lds16 为 1876.38 ms / 143.34 ms / 6.37。因此 O8 将紧凑路径的 TPOT 降低约 8.9%，但它仍不能恢复原布局的带宽。

以下是带 DMA-BUF GPUDirect RDMA、128-token chunk、相同 HTTP workload 的五轮正式矩阵中位数。非紧凑列保留 [N,K] Decode 权重与 [K,N] Prefill 副本；紧凑列释放前者。两次运行使用相同硬件和服务参数，但为独立服务启动，表中应按趋势而不是微小绝对差解释。

场景	非紧凑 TTFT / TPOT / token/s	紧凑 `wave-o8` TTFT / TPOT / token/s
B1, 562 input / 128 output	1885.36 / 108.30 ms / 8.18	1875.64 / 130.52 ms / 6.94
B4 short	918.73 / 95.38 ms / 39.29	920.10 / 182.87 ms / 21.20
B8 short	1704.78 / 123.65 ms / 58.81	1705.54 / 254.81 ms / 30.05
B16 short	3269.73 / 138.84 ms / 97.99	3269.06 / 246.12 ms / 59.31
Context 2048 / 16 output	6789.37 / 156.05 ms / 1.75	6768.73 / 177.82 ms / 1.69
Context 4096 / 16 output	16370.77 / 223.77 ms / 0.81	16353.38 / 245.49 ms / 0.80

结论是明确的：紧凑路径的 Prefill TTFT 几乎不变，但 B1 Decode 仍慢约 20%，而 B4/B8/B16 会回退到 [K,N] hipBLAS 并显著降低吞吐。默认性能部署应保持 --compact-prepacked-gate-up=0；只有显存容量优先、且接受这一回退时才启用紧凑模式。正式工件位于 build-perf/qwen3-32b-tp4pp2/compact-gateup/compact-wave-o8-formal-r5-20260730/ 和 build-perf/qwen3-32b-tp4pp2/compact-gateup/compact-lds16-b1-r5-20260730/。

BF16 AllReduce 的逐行数值特征

TP4 的 hidden-state reduction 使用 ncclBfloat16。RCCL 的 ring/chunk 归约会让不同数据块采用不同的四 rank 加法顺序；而 BF16 加法不是结合律。因此，即使每个 rank 内的多个输入 row 位级相同，不同 row 也可能得到一档 BF16 ULP 内的不同结果。这不是 workspace alias、pipeline transport 覆盖或 rank 间 collective 错序：每个 TP rank 收到的 reduction 输出仍必须逐字节一致。

可用下面的可选四卡诊断复现服务的 Decode M=4,N=5120 与 Prefill M=36/63,N=5120 shape。它对每个 rank 构造一条伪随机 BF16 row，再将该 row 位级复制到所有 batch row，并分别测试 in-place/out-of-place reduction：

METAINFER_REF_ENABLE_QWEN32_RCCL_NUMERICS=1 \
  build-hip/metainfer_ref_rccl_tensor_parallel_collective_test \
  --gtest_filter=RcclTensorParallelCollective.CharacterizesBf16ReductionRowNumericsAtQwen32ServiceShapes

2026-07-28 的结果中，所有 shape 和两种 buffer 模式均为 rank_disagreements=0；跨 row 差异的最大绝对值为 0.0625，相对“FP32 求和后再舍入 BF16”的最大偏差为 0.03125。这正是预期的一档 BF16 舍入差。默认服务不提升为 FP32 AllReduce，因为它会增加转换开销并把该路径的通信数据量从 BF16 的 2 bytes/element 提高到 4 bytes/element；只有需要严格 batch-position bitwise 一致性的专门精度模式才值得采用该折衷。

TP4 x PP2 fenced Rank Timeline

32B 双节点服务的 operator Timeline 是一项诊断工具，不是 TTFT、TPOT 或吞吐的正式测量方式。普通 HIP API 调用只会异步提交 GPU 工作，不能把 host scope 直接当作 kernel 执行范围。诊断运行必须同时开启以下三个变量；此时每个被标记的 GPU operator scope 在退出前同步其所属 stream，允许本机 ROCprof kernel 时间戳按严格包含关系归属到该 scope：

export METAINFER_REF_TIMELINE=1
export METAINFER_REF_TIMELINE_FENCED_DIAGNOSTIC=1
export METAINFER_REF_TIMELINE_OUT=/tmp/qwen32-node0.timeline.jsonl

该模式会改变执行节奏，因此只能用于组成分析。五轮 TTFT/TPOT/throughput 仍必须关闭这三个 Timeline 变量，在相同的常驻 HTTP 服务上单独运行。TCP 控制范围没有 GPU stream，不会进入 ROCprof manifest；QKV、Attention、TP collective、LM Head，以及实际提交 RCCL 工作的 activation/result send/recv 都会产生 fenced GPU scope。

在每个节点分别用本机 rocprofv2 --kernel-trace 完成诊断工作负载并停止服务后，将该节点 JSONL 转成 manifest。--gpu-ids 仅在 ROCprof 报告的 GPU_ID 与 HIP device_ordinal 不相同时需要指定，顺序对应本节点 TP rank 0,1,2,3：

python3 tools/build_tp4pp2_timeline_manifest.py \
  --input /tmp/qwen32-node0.timeline.jsonl \
  --node-rank 0 --workload b4-short \
  --output /tmp/qwen32-node0.manifest.json

python3 tools/build_tp4pp2_timeline_manifest.py \
  --input /tmp/qwen32-node1.timeline.jsonl \
  --node-rank 1 --workload b4-short \
  --output /tmp/qwen32-node1.manifest.json

转换器 fail-fast：每个 GPU scope 都必须同时具有 fence_requested=true、fenced_diagnostic=true，以及正的 roctracer_host_start_ns/end_ns；每节点也必须覆盖固定的四个 Global Rank（Node 0 为 0..3，Node 1 为 4..7）。这可防止普通无 fence 的 host 范围被误用于 kernel 归因。

最后用两个节点各自的 ROCprof CSV 合并；产物只保留每个节点内部的 duration 和 rank imbalance，明确不进行跨节点绝对时间对齐或相减：

python3 tools/run_tp4pp2_kernel_timeline.py \
  --node0-kernel-csv /tmp/node0/results_node0.csv \
  --node0-manifest /tmp/qwen32-node0.manifest.json \
  --node1-kernel-csv /tmp/node1/results_node1.csv \
  --node1-manifest /tmp/qwen32-node1.manifest.json \
  --output build-perf/qwen3-32b-tp4pp2/b4-short-fenced-timeline.json

Name		Name	Last commit message	Last commit date
.codex/skills		.codex/skills
MetaInfer-00 @ a13654c		MetaInfer-00 @ a13654c
benchmarks		benchmarks
cmake		cmake
docs/superpowers		docs/superpowers
include/metainfer_ref		include/metainfer_ref
src		src
testdata/m3		testdata/m3
tests		tests
third_party		third_party
tools		tools
.gitignore		.gitignore
.gitmodules		.gitmodules
CMakeLists.txt		CMakeLists.txt
LANGUAGE_BOUNDARY.md		LANGUAGE_BOUNDARY.md
README.md		README.md
SINGLE_GPU_OPTIMIZATION_PLAYBOOK.md		SINGLE_GPU_OPTIMIZATION_PLAYBOOK.md
TENSOR_PARALLEL_OPTIMIZATION_PLAYBOOK.md		TENSOR_PARALLEL_OPTIMIZATION_PLAYBOOK.md
continuesbatching.md		continuesbatching.md
model_manifest.json		model_manifest.json
model_manifest_qwen3_14b.json		model_manifest_qwen3_14b.json
model_manifest_qwen3_4b.json		model_manifest_qwen3_4b.json
model_manifest_qwen3_8b.json		model_manifest_qwen3_8b.json
pagedkvcache.md		pagedkvcache.md
qwen36_27b_tp4_pp2_optimization.md		qwen36_27b_tp4_pp2_optimization.md
qwen3_14b_tp4_framework.md		qwen3_14b_tp4_framework.md
qwen3_14b_tp4_optimization.md		qwen3_14b_tp4_optimization.md
qwen3_tp4_pp2_framework.md		qwen3_tp4_pp2_framework.md
qwen3_tp4_pp2_optimization.md		qwen3_tp4_pp2_optimization.md
serve.sh		serve.sh
tensor_parallel.md		tensor_parallel.md

Latest commit

History

Folders and files

Repository files navigation

Qwen3 C++ Reference Runtime

可审计性能证据

M0 前置条件

依赖同步

M0 验收

M1 Core 与 Backend

M2 Model Inputs 与 Weight Loading

M3 Golden Oracle（仅生成时）

M3 Numerical Execution 验收

M4a Paged KV 与动态批处理

M4b Chunked Prefill

M5 Sampling 与 OpenAI Service

M6 Reliability 与 Performance Evidence

M6.5 Performance Hardening

M6.5-E Decode Tail 与内存优化

M6.5-F Device Sampling 与 LM Head 候选

Qwen3-4B 单卡 Bring-up

Qwen3-4B 五轮正式性能基线

Qwen3-4B TP2 Bring-up

TP2 五轮性能结果

TP2 HTTP/OpenAI 服务

Qwen3-8B TP2 Bring-up

Rank 0-only LM Head/Sampling

Decode small-M Linear

Decode packed small-M Linear v2

BF16 TP2 性能收尾

Qwen3 TP4 与 RCCL

Qwen3-8B TP4 基线

Qwen3-14B TP4 Bring-up

Qwen3-14B Linear Shape 优化

Qwen3-14B TP4 双 RCCL communicator

Qwen3-14B TP4 B4 multi-output-row GEMV

Qwen3-14B TP4 Persistent Rank Workers

Qwen3-14B TP4 B4 tiled LDS input cache

Qwen3-14B TP4 BF16x2 Paired Decode Attention

Qwen3-14B TP4 Final BF16 Decode 收尾

Qwen3-14B TP4 16x4096 KV Capacity

Qwen3-14B TP4 Performance Matrix

Qwen3-14B TP4 Large-M Prefill 收尾

Tiled Prefill Attention 端到端 A/B

hipBLASLt / Composable Kernel 实验接入

Qwen3-14B TP4 B8/B16 Decode 收尾

B8/B16 LDS Input-Reuse GEMM

Qwen3-14B TP4 W8A16 Decode 与 Selective Prefill

INT8 Paged KV Cache 与 Attention

Qwen3-32B 双节点 TP4 x PP2 短期版本

构建和模型前提

启动顺序

HTTP/OpenAI 验证

当前验收边界

Qwen3-32B Gate/Up 紧凑预打包的容量与性能边界

BF16 AllReduce 的逐行数值特征

TP4 x PP2 fenced Rank Timeline

About

Resources

Stars

Watchers

Forks

Releases

Packages

Contributors

Languages