# LingBot 全流水线分析：首轮报告（持续补充）

生成时间：2026-09-10T17:35:00.408136+08:00

当前结论只覆盖已核验的数据。完整计划尚未完成；所有局部 Attention 结果均不代表整模型优化已经通过。
## 首页摘要：16项验收状态与证据入口

验收表生成时间 2026-09-10T09:23:04.469687+00:00（UTC）；`full_plan_complete`=False。状态只来自 analysis/plan_acceptance_audit.json，本摘要不会把部分完成升级为通过。完成路线见 [completion_plan_20260910.md](/completion_plan_20260910.md)（2026-09-10 用户确认：16项全部通过或有边界闭合；源码统一冻结Case D；数值验收在eager模式对exact参考按既定门槛；H1–H5每个候选一个experiment_id）。

|编号|要求|状态|剩余缺口（摘录）|证据入口|
|---|---|---|---|---|
|K0|环境与完整时间线|部分完成|环境身份及worker尾部覆盖已核对；全服务阶段与CPU因果链仍不完整。 新增同一运行的nsys+NVTX账本（生产模式）与逐chunk关键路径：480x832 稳态DiT 678 ms/VAE段 169 ms（解码 142，IPC回传 16），4卡同时空闲 58 ms；704x…|[original_nsys_worker_flush_v2_anchor_audit.json](/original_nsys_worker_flush_v2_anchor_audit.json) · [worker_longest_gap_context.json](/worker_longest_gap_context.json) · `analysis/ledger_capture_480x832_v1_anchor_audit.json`|
|K1|每次调用、shape和上下文追溯|部分完成|1091749次记录调用有物理GPU/API关联；完整逐层shape对应尚缺。 新增480×832 TP1/SP4的32000次Attention有效长度/阶段/层记录；与原nsys不属同一运行，不能直接逐调用嫁接。 高分辨率新增28个代表输入、90个kernel；两分辨率合计5…|[worker_flush_v2_runtime_link_summary.json](/worker_flush_v2_runtime_link_summary.json) · [worker_launch_resources_summary.json](/worker_launch_resources_summary.json) · [coverage.json](/coverage.json)|
|K2|逐组硬件指标与瓶颈反证|部分完成／已恢复采集|进程级授权与普通用户调度实测兼容；18个GEMM/VAE/转换真实输入计数器及同卡FA2/FA3对照通过审核。完整逐组覆盖、硬件反证及全模型MFU仍缺。 新增13种真实输入/48个kernel已审核；现有31种算子代表输入覆盖此480×832 TP1/SP4逐调用记录的12种线性…|`raw/ncu_attention_fa3_pilot_v1/selected_metrics.json` · [capability_counter_probe_audit.json](/capability_counter_probe_audit.json) · [capability_operator_ncu_full_audit.json](/capability_operator_ncu_full_audit.json)|
|A0|真实Shape×Backend矩阵|部分完成|12配置代表输入已覆盖，1004有效回放、4个math OOM；不是全层覆盖，MFU图尚缺。 新增短KV、长KV及cross四组FA2/FA3共8次NCU均通过数值检查，并按输入SHA256匹配已有未插桩回放；两种计时独立呈现，不推导整模型加速。 新增Attention有效FLO…|[matrix_coverage.json](/matrix_coverage.json) · [fa3_interface_dispatch_comparison.json](/fa3_interface_dispatch_comparison.json) · [attention_hardware_extension_summary_v1.json](/attention_hardware_extension_summary_v1.json)|
|G0|其他重要kernel图谱|部分完成|18个真实输入有NCU、shape和标称FLOPs/DRAM流量图；新增同卡三启动连续访问带宽参考和NCU流量核对；同卡计算27点和12个同卡算子经验参考图已补充；完整模型映射、L2口径及完整Roofline仍缺。 新增13种真实输入/48个kernel已审核；现有31种算子代表…|[operator_replay_v2_audit.json](/operator_replay_v2_audit.json) · [local_gemm_calibration_audit.json](/local_gemm_calibration_audit.json) · [operator_work_traffic.json](/operator_work_traffic.json)|
|G1|整模型优化验收|未通过|cuDNN cross候选40份输出均未通过门槛；尚无接受的整模型加速。 新SDPA Flash候选的6个实际实现输入已完成独立数值和kernel审核；该结果不能替代整模型验收。 本次检查未确认模型任务仍存活，不能仅凭历史状态认定在运行。 新候选保存输出数值检查未通过；单次A/B…|[cross_candidate_eager_r0_comparison.json](/cross_candidate_eager_r0_comparison.json) · [execution_control_pilot_timings.json](/execution_control_pilot_timings.json) · [cross_sdpa_flash_impl_audit.json](/cross_sdpa_flash_impl_audit.json)|
|H0|参考优化适用性|部分完成|已做本地源码映射；所列外部PR/RFC状态及本地提交祖先关系已追加核验，完整编译融合及transport运行关联仍缺。|[optimization_transfer_matrix.json](/optimization_transfer_matrix.json) · [h3_external_reference_review.json](/h3_external_reference_review.json) · [h3_reference_local_ancestry.json](/h3_reference_local_ancestry.json)|
|H1|Attention与通信布局消融|未完成|独立后端回放不能替代通信布局的单变量A/B。|[optimization_transfer_matrix.json](/optimization_transfer_matrix.json)|
|H2|DiT融合与局部数据构造消融|未完成|尚无通过数值门槛的融合增量收益验证。|[optimization_transfer_matrix.json](/optimization_transfer_matrix.json)|
|H3|解码算子与空间并行消融|未完成|新增含全部转换成本的局部布局A/B：数值逐值一致但慢约22.42%；仅单输入局部反证，完整解码优化消融仍缺。 新增时间维卷积全成本布局A/B同样逐值一致但慢14.82%；布局转换51.9%时间占比未直接转化为收益。|[optimization_transfer_matrix.json](/optimization_transfer_matrix.json) · [missing_vae_components_v1.json](/missing_vae_components_v1.json) · [vae_layout_temporal_ab_audit_v1.json](/vae_layout_temporal_ab_audit_v1.json)|
|H4|输出与传输优化|未完成|已有uint8路径与交接计时，尚无受控输出优化及编码后媒体验收。 已补充六阶段生产者/消费者表：冻结路径只回传元数据，不含像素IPC；CPU-ready不能替代编码或客户端交付。输出融合真实输入补采已提交，尚未验收。 新增固定FFV1无损CPU编码往返，3份视频和351帧逐字节审…|[clean_handoff_summary.json](/clean_handoff_summary.json) · [optimization_transfer_matrix.json](/optimization_transfer_matrix.json) · [output_boundary_inventory_v1.json](/output_boundary_inventory_v1.json)|
|H5|流水线重叠、背压与取消|未完成|时间线重叠不等于优化收益；慢消费者、取消和组合调度验收尚缺。|[worker_chunk_activity.json](/worker_chunk_activity.json) · [optimization_transfer_matrix.json](/optimization_transfer_matrix.json)|
|TAIL|重复运行与尾延迟|部分完成|两分辨率各3次、每次1000间隔；新增连续会话块重采样显示部分P99区间敏感，数值控制与未来稳定性仍未完成。|`raw/tail_timing_audit_workflow_v1.json` · [clean_handoff_summary.json](/clean_handoff_summary.json) · [tail_six_launch_hash_repeatability.json](/tail_six_launch_hash_repeatability.json)|
|NUM|数值可重复性与候选门槛|部分完成|eager两次40份输出逐位一致；两种编译模式重复及cuDNN候选均失败，原因未最终定位。 新参考与历史40份逐值一致；原版和候选两条轨迹各160次同输入比较及FP64抽样通过，诊断最终输出各40份与对应未插桩运行逐值一致。跨轨迹Q从第2次出现差异，仍缺完整逐层因果证明。 新增分…|[execution_control_eager_repeat_comparison.json](/execution_control_eager_repeat_comparison.json) · [execution_error_progression.json](/execution_error_progression.json) · [cross_sdpa_fresh_reference_repeat_v2.json](/cross_sdpa_fresh_reference_repeat_v2.json)|
|RAW|全量原始数据与复现|部分完成|已存在文件有索引、源快照与复算脚本；长尾后续轮次未保存完整输出，无法补回。|[delivery_archive_audit_20260910T0020Z.json](/delivery_archive_audit_20260910T0020Z.json) · [REPRODUCE_CN.md](/REPRODUCE_CN.md) · [raw_data_index.json](/raw_data_index.json)|
|WEB|首轮报告与证据访问|首版已交付／持续更新|中文报告及看板已发布，82个报告内资源链接通过内容核对；不代表完整计划完成。|[REPORT_CN.md](/REPORT_CN.md) · [public_report_asset_audit.json](/public_report_asset_audit.json)|

已知负结果（保留，不重试同口径）：

- 两种编译模式（default、reduce-overhead）跨启动重复均 0/40 通过门槛；eager 跨启动逐位一致，是候选判定的唯一参考。
- cuDNN cross 与 SDPA Flash cross 候选：局部通过，整模型 0/40 通过门槛；不接受。
- 两个含全部转换成本的 VAE channels-last 局部 A/B：逐值一致但分别慢 22.42%、14.82%；不接受。
- 长尾 6 次运行仅 epoch0/1 保存完整输出，其余不可补回；最终验收长尾将重跑并保存全部输出。
- 704×1280 稳态约 4.89 FPS，与 12 FPS 目标差约 2.46 倍；480×832 约 14.4 FPS 已高于 12。

唯一局部正结果：输出转换 Triton 融合，20 份真实输入与 775 个取整边界值逐字节一致，含 D2H 局部耗时降低 23.76%/31.84%；尚未集成整模型验收。

## 当前可用结论

首轮固定报告包已在2026-09-10 08:40（北京时间）交付，以下为持续补充的在线版本；完整研究计划仍未全部完成。

- Attention：12个配置的代表上下文共尝试1008个回放点，1004个通过审核，4个SDPA math点显存不足；并非所有模型层完整覆盖。
- 硬件证据：59种真实线性/卷积/转换代表输入共183次NCU kernel记录（含一项调度异常备注），以及同卡FA2/FA3对照已审核；原始输入、计数器和输出证据可复用。
- 同卡参考：计算与连续读写带宽校准已完成，12个算子纳入经验参考图；这不是整模型MFU或理论硬件峰值。
- 优化验收：cuDNN cross候选未通过整模型误差门槛；含全部转换成本的VAE局部channels-last候选反而慢约22.42%，不认领加速。SDPA Flash cross整模型候选40份输出也全部未通过门槛；新参考40份逐值一致，失败证据与单次耗时见末节。
- 性能边界：现有交付计时主要到CPU-ready uint8；编码、网络、客户端播放、背压与取消仍未建立完整证据。
- 数据限制：六次长尾运行保存了逐次计时/哈希，但后续轮次未保留完整输出；已缺失的完整输出不能从哈希恢复。所有新实验继续保存独立目录和校验清单。

后文保留历史失败、诊断与各次补充，均按测量范围解释；旧阶段的限制如已解决，以对应更新章节为准。


**首轮固定版本报告包**：[下载tar.gz（10.54MB）](/lingbot_first_report_20260910T004008Z.tar.gz) · [SHA256校验文件](/lingbot_first_report_20260910T004008Z.tar.gz.sha256)。包含报告、图表、下载数据、脚本和索引；不包含约326GB完整原始实验文件。解压后的README_CN.md说明离线查看方式和原始数据位置。在线报告继续更新，固定包保留打包时点内容。

## 首轮报告：当前可确认的结论

以下结论优先于后文按实验阶段保留的历史记录。完整计划尚未完成。

|问题|当前证据与结论|
|---|---|
|Attention输入与回放|两种分辨率、六种TP/SP组合的纯DiT输入已覆盖；1008项回放中1004项有效，4项SDPA math显存不足。孤立回放不能代替整模型验收。|
|正常运行速度|六次扩样运行已完成。480×832的CPU-ready吞吐约14.38–14.55 FPS；704×1280约4.87–4.89 FPS。每次1000个稳态交付间隔，不包含编码、网络和客户端。|
|数值可重复性|新增无诊断钩子对照中，eager两次启动的40份输出逐位一致；reduce-overhead与default编译模式各自重复比较均0/40通过数值门槛。仅适用于已测配置；未证明唯一原因。|
|cuDNN cross候选|eager整模型候选与可重复的eager参考比较，40份输出均未通过预设门槛；不接受为有效优化。选择日志已保存，本次候选没有新增kernel时间线。|
|kernel硬件分析|已有真实FA3输入的NCU试采。矩阵乘、VAE卷积和转换kernel的完整硬件指标仍缺失；root采集被调度守护终止后未绕过限制。不能给出完整MFU或Roofline结论。|
|数据保存边界|原始文件、失败日志、时间线、张量和校验索引位于持久目录。长尾实验只有前两轮保留完整输出，后续轮次保留计时与哈希，无法由哈希还原输出；不宣称所有轮次完整输出已归档。|

新增七次短运行每次仅四个稳态计时间隔，不能作尾延迟或稳定性证明。计时复算脚本：scripts/reproduce_execution_control_pilot_timings.py；验证结果：analysis/execution_control_pilot_timing_reproduction.txt。详细数值、范围和证据路径见下文。

## 原计划逐项验收状态

首轮报告已交付，完整研究计划仍未通过验收。以下保留原计划全部工作包，不将已完成的微基准或时间线核对替代全流水线归因。

|工作包|要求|状态|证据能支持的范围及缺口|
|---|---|---|---|
|K0|环境与完整时间线|部分完成|环境身份及worker尾部覆盖已核对；全服务阶段与CPU因果链仍不完整。 新增同一运行的nsys+NVTX账本（生产模式）与逐chunk关键路径：480x832 稳态DiT 678 ms/VAE段 169 ms（解码 142，IPC回传 16），4卡同时空闲 58 ms；704x1280 稳态DiT 2138 ms/VAE段 338 ms（解码 290，IPC回传 34），4卡同时空闲 59 ms。时钟锚点与尾部检查通过。仍缺：CPU侧等待的逐线程因果链（请求进入与结果返回的空闲段仅定位到进程边界）与编码/网络/客户端段。|
|K1|每次调用、shape和上下文追溯|部分完成|1091749次记录调用有物理GPU/API关联；完整逐层shape对应尚缺。 新增480×832 TP1/SP4的32000次Attention有效长度/阶段/层记录；与原nsys不属同一运行，不能直接逐调用嫁接。 高分辨率新增28个代表输入、90个kernel；两分辨率合计59个算子输入覆盖各自所选线性/卷积/转换结构。PID2019660有未归因的调度备注，保留原失败和后续带备注审核。两分辨率逐调用元数据均已补充，仍非与旧nsys同次运行映射。 新增同一运行逐调用账本：480x832 1,092,983次kernel，87.6%调用/99.9%稳态GPU时间有NVTX上下文，graph内位置由inductor wrapper源码对齐到子算子；704x1280 1,092,185次kernel，87.5%调用/100.0%稳态GPU时间有NVTX上下文，graph内位置由inductor wrapper源码对齐到子算子。shape按eager全算子dispatch采集附加（480x832 稳态GPU时间94.7%有shape；704x1280 稳态GPU时间95.8%有shape），融合kernel的shape为推断值。仍缺：融合kernel精确字节、非同一运行的shape来源说明已保留。 eager全算子dispatch采集的40份输出与对应eager参考逐位一致：480x832、704x1280。|
|K2|逐组硬件指标与瓶颈反证|部分完成／已恢复采集|进程级授权与普通用户调度实测兼容；18个GEMM/VAE/转换真实输入计数器及同卡FA2/FA3对照通过审核。完整逐组覆盖、硬件反证及全模型MFU仍缺。 新增13种真实输入/48个kernel已审核；现有31种算子代表输入覆盖此480×832 TP1/SP4逐调用记录的12种线性、17种卷积与2种转换结构。仍不代表其他配置、所有实际调用或全模型MFU。 高分辨率新增28个代表输入、90个kernel；两分辨率合计59个算子输入覆盖各自所选线性/卷积/转换结构。PID2019660有未归因的调度备注，保留原失败和后续带备注审核。两分辨率逐调用元数据均已补充，仍非与旧nsys同次运行映射。 新增短KV、长KV及cross四组FA2/FA3共8次NCU均通过数值检查，并按输入SHA256匹配已有未插桩回放；两种计时独立呈现，不推导整模型加速。 新增全模型标称FLOPs与模型级MFU（同一运行关键路径为分母）：480x832 DiT 49.0%（BF16 dense）、VAE 14.5%（TF32 dense）、端到端 40.2%；704x1280 DiT 53.1%（BF16 dense）、VAE 15.9%（TF32 dense）、端到端 46.6%。标称口径，非硬件执行FLOPs。 新增69张逐组瓶颈卡（unresolved_no_dominant_limiter 17, compute_bound 29, partial_wave_tail 13, dram_bandwidth_bound 4, l2_bandwidth_bound 6）与账本分组归因：480x832 计数器归因68%/通信12%/推导6%/未归因14%；704x1280 计数器归因78%/通信7%/推导4%/未归因11%。稳态KV长度FA3核心按阈值规则无主导单元，stall/tensor-pipe/L2重采已排队；compiled融合kernel与paged gather无计数器；全模型MFU仍缺。|
|A0|真实Shape×Backend矩阵|部分完成|12配置代表输入已覆盖，1004有效回放、4个math OOM；不是全层覆盖，MFU图尚缺。 新增短KV、长KV及cross四组FA2/FA3共8次NCU均通过数值检查，并按输入SHA256匹配已有未插桩回放；两种计时独立呈现，不推导整模型加速。 新增Attention有效FLOP利用率曲线：1004个聚合点，分母为规格BF16 dense 989.5 TF/s，另标同卡经验上限；回放GPU未逐卡校准，不是整模型MFU。|
|G0|其他重要kernel图谱|部分完成|18个真实输入有NCU、shape和标称FLOPs/DRAM流量图；新增同卡三启动连续访问带宽参考和NCU流量核对；同卡计算27点和12个同卡算子经验参考图已补充；完整模型映射、L2口径及完整Roofline仍缺。 新增13种真实输入/48个kernel已审核；现有31种算子代表输入覆盖此480×832 TP1/SP4逐调用记录的12种线性、17种卷积与2种转换结构。仍不代表其他配置、所有实际调用或全模型MFU。 小工作集三启动Graph配对和NCU流量已审核；Graph平均调用开销与普通提交显著不同，存在DRAM流量，仍不能认领纯L2峰值。 高分辨率新增28个代表输入、90个kernel；两分辨率合计59个算子输入覆盖各自所选线性/卷积/转换结构。PID2019660有未归因的调度备注，保留原失败和后续带备注审核。两分辨率逐调用元数据均已补充，仍非与旧nsys同次运行映射。 新增全模型标称FLOPs与模型级MFU（同一运行关键路径为分母）：480x832 DiT 49.0%（BF16 dense）、VAE 14.5%（TF32 dense）、端到端 40.2%；704x1280 DiT 53.1%（BF16 dense）、VAE 15.9%（TF32 dense）、端到端 46.6%。标称口径，非硬件执行FLOPs。 新增DRAM层Roofline（67个点，含Attention NCU）；L2层计数器未采集，重采已排队；完整模型映射改由同一运行账本承担。|
|G1|整模型优化验收|未通过|cuDNN cross候选40份输出均未通过门槛；尚无接受的整模型加速。 新SDPA Flash候选的6个实际实现输入已完成独立数值和kernel审核；该结果不能替代整模型验收。 本次检查未确认模型任务仍存活，不能仅凭历史状态认定在运行。 新候选保存输出数值检查未通过；单次A/B不构成稳定性能验收。|
|H0|参考优化适用性|部分完成|已做本地源码映射；所列外部PR/RFC状态及本地提交祖先关系已追加核验，完整编译融合及transport运行关联仍缺。|
|H1|Attention与通信布局消融|未完成|独立后端回放不能替代通信布局的单变量A/B。|
|H2|DiT融合与局部数据构造消融|未完成|尚无通过数值门槛的融合增量收益验证。|
|H3|解码算子与空间并行消融|未完成|新增含全部转换成本的局部布局A/B：数值逐值一致但慢约22.42%；仅单输入局部反证，完整解码优化消融仍缺。 新增时间维卷积全成本布局A/B同样逐值一致但慢14.82%；布局转换51.9%时间占比未直接转化为收益。|
|H4|输出与传输优化|未完成|已有uint8路径与交接计时，尚无受控输出优化及编码后媒体验收。 已补充六阶段生产者/消费者表：冻结路径只回传元数据，不含像素IPC；CPU-ready不能替代编码或客户端交付。输出融合真实输入补采已提交，尚未验收。 新增固定FFV1无损CPU编码往返，3份视频和351帧逐字节审核通过；仍非模型服务集成、像素IPC或有损质量验收。 交接终态：输出补采与局部融合A/B已完成；40模型输出、20输入转换与775边界字节一致，D2H-inclusive局部耗时降低23.76%/31.84%；未集成模型。用户要求停止新增实验并交接。|
|H5|流水线重叠、背压与取消|未完成|时间线重叠不等于优化收益；慢消费者、取消和组合调度验收尚缺。|
|TAIL|重复运行与尾延迟|部分完成|两分辨率各3次、每次1000间隔；新增连续会话块重采样显示部分P99区间敏感，数值控制与未来稳定性仍未完成。|
|NUM|数值可重复性与候选门槛|部分完成|eager两次40份输出逐位一致；两种编译模式重复及cuDNN候选均失败，原因未最终定位。 新参考与历史40份逐值一致；原版和候选两条轨迹各160次同输入比较及FP64抽样通过，诊断最终输出各40份与对应未插桩运行逐值一致。跨轨迹Q从第2次出现差异，仍缺完整逐层因果证明。 新增分步诊断：default编译模式（无CUDA Graph）两次启动0/40通过；reduce-overhead加库级确定性控制（cuBLAS工作区、cuDNN确定性、use_deterministic_algorithms）两次启动40/40通过、40/40逐位一致：跨启动差异定位为库级非确定性而非编译/Graph。单一控制拆分与性能代价计时排队中。|
|RAW|全量原始数据与复现|部分完成|已存在文件有索引、源快照与复算脚本；长尾后续轮次未保存完整输出，无法补回。|
|WEB|首轮报告与证据访问|首版已交付／持续更新|中文报告及看板已发布，82个报告内资源链接通过内容核对；不代表完整计划完成。|

[下载逐项验收表及证据SHA256](/plan_acceptance_audit.json)。此表是审核时点快照；引用文件之后可能继续更新，快照中的哈希用于识别当时版本。首版报告交付不等于所有要求通过。

## 历史记录：旧时间线窗口与尾部限制

此前整模型诊断记录中，最后一个具名FA3 kernel早于按当前公式映射的最后DiT结束时间数秒。此现象尚不能证明时钟公式错误，也可能涉及CUDA Graph采集覆盖或日志边界含其他工作。这些旧记录的measured/steady窗口归属不用于完整阶段性能结论。完整时间线中的具名kernel存在性不依赖此窗口转换。四份时间线均已检查进程末尾：部分DiT记录提前结束，退出日志存在worker超时终止记录。旧时间线的稳态kernel归属也暂列待核验；已记录事件数不代表真实调用覆盖100%。这不改变独立CPU-ready计时数据。后续定时写出补采仍缺少部分尾部；再后的worker显式写出v2已通过独立时钟锚点和各worker末块存在性检查，详见下方v2结果。该通过结果不修复本节旧记录。证据：analysis/integrated_window_consistency_review.json、analysis/trace_tail_audit_v1.json；补采任务：raw/flush_trace_workflow_v1.json。

## 本机计算校准任务

状态：completed_pending_audit。在时间线补采之后，普通用户通过gpu run申请一张卡，运行三次独立进程；各测4096、8192、12288方阵的BF16、FP32禁用TF32和FP32允许TF32配置。每项20次预热、100次CUDA event样本，实际kernel另行记录。审核结果见单卡大矩阵计算校准实测一节；这只提供单卡大矩阵的实测参考，不是硬件规格峰值，也不是模型MFU。证据：raw/calibration_workflow_v1.json。

## 单卡大矩阵计算校准实测

审核状态：passed。三次独立进程，每个配置20次预热、100次CUDA event样本。表中为三次运行各自中位时长换算的TFLOP/s范围，2MNK为算法FLOPs。实际kernel另行采集，未混入计时。

|输入/计算策略|M=N=K|三次实测TFLOP/s范围|
|---|---:|---:|
|bf16|4096|787.06–788.28|
|bf16|8192|693.82–694.28|
|bf16|12288|706.99–707.85|
|fp32_ieee|4096|51.17–51.20|
|fp32_ieee|8192|51.25–51.29|
|fp32_ieee|12288|50.32–50.34|
|fp32_tf32_allowed|4096|394.31–395.98|
|fp32_tf32_allowed|8192|360.68–363.50|
|fp32_tf32_allowed|12288|357.65–358.46|

物理设备：_CudaDeviceProperties(name='NVIDIA L20X', major=9, minor=0, total_memory=143166MB, multi_processor_count=132, uuid=006eb78c-23cb-f37d-eb7c-0ccb578b8f11, pci_bus_id=8, pci_device_id=0, pci_domain_id=0, L2_cache_size=60MB)。仅校准这张卡，不将其当作所有GPU或所有shape的上限。允许TF32与禁用TF32必须分开解释；BF16独立Attention的有效FLOP吞吐也不能直接等同大GEMM利用率。尚未完成硬件规格上限核验及DRAM/L2字节测量，不生成完整roofline或模型MFU结论。[下载校准CSV与完整kernel名称](/local_gemm_calibration.csv)。证据：analysis/local_gemm_calibration_audit.json。

## 未插桩基线扩样安排

状态：六次运行均已结束并通过计时审核。在前序采集/校准之后，两种分辨率各三次独立启动，每次1轮预热+250个reset session，目标1000个稳态交付间隔。基于原始参考实现，不表示rebase或cuDNN候选已通过数值验收。受GPU等待和08:30截止限制，未完成时保留实际样本数，不宣称大样本验收通过。状态与完整命令：raw/tail_baseline_workflow_v1.json。

## 480×832三次独立运行的差异

三次均完成250个测量session、1000个稳态间隔，并通过独立时间戳复核。FPS范围14.379–14.550，最高相对最低高1.19%；P99范围837.774–865.680 ms，相差3.33%。每次最大间隔范围881.423–1028.816 ms。

这里满足本分辨率三次独立启动、每次1000稳态样本的数量要求；范围仅描述这三次观测，没有合并样本缩小置信区间，也不能将观测最大值作为未来延迟上限。首轮GPU0–3，后两轮GPU4–7；启动状态与设备影响尚未隔离。输出跨启动数值门槛未通过，不能据此接受优化候选。高分辨率重复实验另行报告。

证据：analysis/tail_480_three_launch_summary.json，包含各次metrics文件SHA256。

## 704×1280三次独立运行汇总

三次均完成250个测量session、1000个稳态交付间隔，并通过原始时间戳独立复核。均使用GPU4–7；CPU-ready uint8边界，不含编码和网络。

|运行|FPS|P99 ms|最大间隔 ms|
|---|---:|---:|---:|
|tail_original_704x1280_r0_v1|4.888594|2597.353|2899.883|
|tail_original_704x1280_r1_v1|4.872355|2664.672|2980.876|
|tail_original_704x1280_r2_v1|4.888763|2475.593|2711.041|

这些范围只描述三次启动，没有预设稳定性通过阈值，也未合并样本缩小置信区间。观测最大值不是未来上限。保存输出跨启动未通过原数值门槛，不接受优化候选。退出清理阶段有worker及共享资源警告，完整日志保留。

[第三轮逐次计时CSV](/tail_original_704x1280_r2_v1_timings.csv) · [第三轮分布图](/tail_original_704x1280_r2_v1_delivery_distribution.png)。证据：analysis/tail_704_three_launch_summary.json，含各次metrics校验和。

## 大样本基线：已审核结果

CPU审核状态：terminal_runs_audited；目前通过完整性审核的独立运行数：6。仅将成功结束且审核通过的运行列入下表；运行中和失败任务不计为有效完成。

|运行|session数|稳态间隔数|稳态FPS|P50 ms|P95 ms|P99 ms|最大间隔ms|
|---|---:|---:|---:|---:|---:|---:|---:|
|tail_original_480x832_r0_v1|250|1000|14.550|823.968|831.784|837.774|1028.816|
|tail_original_480x832_r1_v1|250|1000|14.379|832.925|845.024|865.680|1010.481|
|tail_original_480x832_r2_v1|250|1000|14.412|832.049|840.781|847.602|881.423|
|tail_original_704x1280_r0_v1|250|1000|4.889|2450.489|2466.998|2597.353|2899.883|
|tail_original_704x1280_r1_v1|250|1000|4.872|2456.933|2476.479|2664.672|2980.876|
|tail_original_704x1280_r2_v1|250|1000|4.889|2453.967|2466.622|2475.593|2711.041|

![全部稳态间隔及累计分布](/tail_original_480x832_r0_v1_delivery_distribution.png)

图保留全部稳态样本和最大停顿，不进行尾部裁剪；仅代表本次独立启动。

tail_original_480x832_r0_v1：P99的进程内session分组bootstrap 95%区间为836.430–842.137 ms；不是独立启动之间的置信区间。

[下载tail_original_480x832_r0_v1全部逐块计时CSV](/tail_original_480x832_r0_v1_timings.csv)。CSV保留非稳态块并提供steady标记；最终稳态统计仅使用每个测量session的间隔5–8，原始时间戳另行重算审核通过。

![全部稳态间隔及累计分布](/tail_original_480x832_r1_v1_delivery_distribution.png)

图保留全部稳态样本和最大停顿，不进行尾部裁剪；仅代表本次独立启动。

tail_original_480x832_r1_v1：P99的进程内session分组bootstrap 95%区间为853.742–885.880 ms；不是独立启动之间的置信区间。

[下载tail_original_480x832_r1_v1全部逐块计时CSV](/tail_original_480x832_r1_v1_timings.csv)。CSV保留非稳态块并提供steady标记；最终稳态统计仅使用每个测量session的间隔5–8，原始时间戳另行重算审核通过。

![全部稳态间隔及累计分布](/tail_original_480x832_r2_v1_delivery_distribution.png)

图保留全部稳态样本和最大停顿，不进行尾部裁剪；仅代表本次独立启动。

tail_original_480x832_r2_v1：P99的进程内session分组bootstrap 95%区间为845.721–861.729 ms；不是独立启动之间的置信区间。

[下载tail_original_480x832_r2_v1全部逐块计时CSV](/tail_original_480x832_r2_v1_timings.csv)。CSV保留非稳态块并提供steady标记；最终稳态统计仅使用每个测量session的间隔5–8，原始时间戳另行重算审核通过。

![全部稳态间隔及累计分布](/tail_original_704x1280_r0_v1_delivery_distribution.png)

图保留全部稳态样本和最大停顿，不进行尾部裁剪；仅代表本次独立启动。

tail_original_704x1280_r0_v1：P99的进程内session分组bootstrap 95%区间为2497.196–2703.041 ms；不是独立启动之间的置信区间。

[下载tail_original_704x1280_r0_v1全部逐块计时CSV](/tail_original_704x1280_r0_v1_timings.csv)。CSV保留非稳态块并提供steady标记；最终稳态统计仅使用每个测量session的间隔5–8，原始时间戳另行重算审核通过。

![全部稳态间隔及累计分布](/tail_original_704x1280_r1_v1_delivery_distribution.png)

图保留全部稳态样本和最大停顿，不进行尾部裁剪；仅代表本次独立启动。

tail_original_704x1280_r1_v1：P99的进程内session分组bootstrap 95%区间为2513.651–2753.144 ms；不是独立启动之间的置信区间。

[下载tail_original_704x1280_r1_v1全部逐块计时CSV](/tail_original_704x1280_r1_v1_timings.csv)。CSV保留非稳态块并提供steady标记；最终稳态统计仅使用每个测量session的间隔5–8，原始时间戳另行重算审核通过。

![全部稳态间隔及累计分布](/tail_original_704x1280_r2_v1_delivery_distribution.png)

图保留全部稳态样本和最大停顿，不进行尾部裁剪；仅代表本次独立启动。

tail_original_704x1280_r2_v1：P99的进程内session分组bootstrap 95%区间为2471.552–2481.631 ms；不是独立启动之间的置信区间。

[下载tail_original_704x1280_r2_v1全部逐块计时CSV](/tail_original_704x1280_r2_v1_timings.csv)。CSV保留非稳态块并提供steady标记；最终稳态统计仅使用每个测量session的间隔5–8，原始时间戳另行重算审核通过。

每次目标250个reset session、1000个稳态间隔；它们来自同一模型进程，不等于1000次独立启动。原始每块计时及session整组bootstrap区间保存在每个analysis/*_timing目录，重复运行之间仍需分别比较。P99不代表客户端SLA，边界仍为CPU-ready uint8。证据：raw/tail_timing_audit_workflow_v1.json。

## 相同GPU组的跨启动输出控制

r1与r2均分配GPU4–7，源码、冻结脚本、输入哈希、计时边界和profiler模式一致。保存的epoch0/1共40项对应输出全部完成比较，均未通过原数值门槛。该结果说明跨启动差异不能仅用GPU组不同解释；仍不能确定启动状态、随机数或算法选择中的具体原因。输出比较与计时审核是独立检查；三轮计时结果见扩样表，计时通过不表示输出数值门槛通过。

|保存阶段|输出|指标|最小值|最大值|通过数|
|---|---|---|---:|---:|---:|
|warmup|latent|relative_l2|0.015967|0.196514|0/10|
|warmup|pixels|rmse|0.010162|0.135849|0/10|
|steady|latent|relative_l2|0.015967|0.196514|0/10|
|steady|pixels|rmse|0.010162|0.135849|0/10|

证据：analysis/tail_480_r1_r2_saved_output_comparison.json（逐文件哈希与误差）、analysis/tail_480_r1_r2_output_summary.json（身份检查）。这些是同版本控制，不是后端优化收益或感知质量认证。

## 704×1280第二、三次运行的保存输出对比

已比较40个保存张量，40个未通过冻结数值门槛。范围仅为epoch 0、1的latent与像素输出，不能推广到全部250个测量session。

- latent相对L2范围：0.029937–0.403371。
- 归一化像素RMSE范围：0.017613–0.291942。

这是原始模型跨启动控制，数值差异原因尚未定位；不能据此接受优化或放宽正确性门槛。逐文件误差与哈希保存在analysis/tail_704_r1_r2_saved_output_comparison.json。

## 704×1280两次原始模型运行的保存输出对比

已比较40个保存张量，40个未通过冻结数值门槛。范围仅为epoch 0、1的latent与像素输出，不能推广到全部250个测量session。

- latent相对L2范围：0.036485–0.309353。
- 归一化像素RMSE范围：0.019445–0.212082。

这是原始模型跨启动控制，数值差异原因尚未定位；不能据此接受优化或放宽正确性门槛。逐文件误差与哈希保存在analysis/tail_704_r0_r1_saved_output_comparison.json。

## 第二轮长尾样本的时间戳分解

第二轮1000个稳态间隔也通过非负与分段求和检查。最大样本位于epoch=32、chunk=8，交付间隔1010.481 ms。其中latent-ready到发出解码为180.323 ms，该段全样本中位数为4.817 ms；该样本DiT区间675.661 ms，decoder计算区间130.184 ms。

本轮最大样本的额外间隔与首轮所在位置不同，不能将所有长尾归因于同一个阶段。这里仅定位主机标记之间的间隔，没有CPU采样和逐事件关联，不能确认具体函数、调度抢占或通信原因。

[下载第二轮1000个间隔分解CSV](/tail_original_480x832_r1_v1_components.csv)。证据：analysis/tail_original_480x832_r1_v1_components/summary.json。

## 第三轮长尾样本的时间戳分解

第三轮1000个稳态间隔通过非负与分段求和检查。最大样本位于epoch=162、chunk=6，交付间隔881.423 ms；DiT主机区间687.571 ms，decoder计算区间131.408 ms。DiT结束到latent-ready为18.194 ms，该段中位数为0.115 ms。

这些数值是同一节点主机时间标记之间的区间，不能证明GPU kernel或CPU调度的具体原因，也不能把不同样本各段最大值相加。

[下载第三轮1000个间隔分解CSV](/tail_original_480x832_r2_v1_components.csv)。证据：analysis/tail_original_480x832_r2_v1_components/summary.json。

## 704×1280第三轮长尾时间戳分解

1000个稳态间隔均通过非负与分段求和检查。最大交付间隔位于epoch=163、chunk=5，总长2711.041 ms。

|主机时间戳区间|全样本中位数ms|最大交付样本中的ms|
|---|---:|---:|
|上一块CPU-ready到主机收到|27.534|27.327|
|主机收到上一块到下一块DiT开始|3.024|3.006|
|DiT主机记录区间|2114.310|2115.700|
|DiT结束到latent-ready|0.136|0.453|
|latent-ready到发出解码|6.849|262.391|
|发出到decoder开始|3.092|3.416|
|decoder传输区间|5.271|5.156|
|decoder计算区间|285.429|283.800|
|decode结束到CPU-ready|3.010|9.792|

该样本latent-ready到发出解码为262.391 ms，该段中位数6.849 ms；DiT与decoder区间接近各自中位数。没有CPU采样证据，不能确定具体原因。这是主机标记区间分解，不是逐kernel硬件归因；各段中位数不能相加作为总间隔中位数。

[下载高分辨率第三轮1000个间隔分解CSV](/tail_original_704x1280_r2_v1_components.csv)。证据：analysis/tail_original_704x1280_r2_v1_components/summary.json。

## 704×1280第二轮长尾时间戳分解

1000个稳态间隔均通过非负与分段求和检查。最大交付间隔位于epoch=11、chunk=7，总长2980.876 ms。

|主机时间戳区间|全样本中位数ms|最大交付样本中的ms|
|---|---:|---:|
|上一块CPU-ready到主机收到|27.570|29.249|
|主机收到上一块到下一块DiT开始|2.994|81.960|
|DiT主机记录区间|2114.988|2301.460|
|DiT结束到latent-ready|0.127|0.189|
|latent-ready到发出解码|6.588|174.909|
|发出到decoder开始|2.841|9.014|
|decoder传输区间|5.306|26.540|
|decoder计算区间|285.785|354.159|
|decode结束到CPU-ready|9.989|3.397|

该样本在DiT区间、发出解码前的间隔及decoder区间均有增加，不能把长尾归于单一阶段。这是主机标记区间分解，不是逐kernel硬件归因；各段中位数不能相加作为总间隔中位数。

[下载高分辨率第二轮1000个间隔分解CSV](/tail_original_704x1280_r1_v1_components.csv)。证据：analysis/tail_original_704x1280_r1_v1_components/summary.json。

## 704×1280首轮长尾时间戳分解

1000个稳态间隔均通过非负与分段求和检查。最大交付间隔位于epoch=217、chunk=6，总长2899.883 ms。

|主机时间戳区间|全样本中位数ms|最大交付样本中的ms|
|---|---:|---:|
|上一块CPU-ready到主机收到|27.576|31.474|
|主机收到上一块到下一块DiT开始|3.025|88.415|
|DiT主机记录区间|2114.965|2290.361|
|DiT结束到latent-ready|0.131|0.079|
|latent-ready到发出解码|6.873|113.945|
|发出到decoder开始|2.840|9.656|
|decoder传输区间|5.297|17.121|
|decoder计算区间|285.858|344.874|
|decode结束到CPU-ready|2.483|3.960|

该样本在DiT区间、发出解码前的间隔及decoder区间均有增加，不能把长尾归于单一阶段。这是主机标记区间分解，不是逐kernel硬件归因；各段中位数不能相加作为总间隔中位数。

[下载高分辨率首轮1000个间隔分解CSV](/tail_original_704x1280_r0_v1_components.csv)。证据：analysis/tail_original_704x1280_r0_v1_components/summary.json。

## 首轮长尾样本的时间戳分解

对1000个稳态交付间隔使用连续主机时间戳分解，各段均非负且相加等于交付间隔。最大样本位于epoch=129、chunk=7，总长1028.816 ms。各段中位数来自不同样本，不能相加当作整体中位数。

|时间戳间隔|1000样本中位数ms|最大交付样本中的ms|
|---|---:|---:|
|上一块CPU-ready到主机收到|12.753|12.517|
|主机收到上一块到下一块DiT开始|2.902|199.510|
|DiT主机记录区间|664.087|670.074|
|DiT结束到latent-ready|0.119|0.234|
|latent-ready到发出解码|4.754|5.330|
|发出到decoder开始|1.677|1.271|
|decoder传输区间|3.379|3.719|
|decoder计算区间|132.198|134.991|
|decode结束到CPU-ready|1.165|1.170|

最大样本的主要额外间隔位于主机收到上一块之后、下一块DiT开始之前（199.510 ms，通常中位数2.902 ms）；DiT为670.074 ms、decoder计算为134.991 ms，接近各自通常水平。这定位到需要进一步观察的控制流程间隔，但没有CPU采样或调度证据，不能归因为某个Python函数、抢占或GPU kernel。

[下载1000个间隔的完整分解CSV](/tail_original_480x832_r0_v1_components.csv)。原始证据及检查：analysis/tail_original_480x832_r0_v1_components/summary.json、raw/tail_original_480x832_r0_v1/out/summary.json。

## 扩样两次启动的保存输出对照

r0与r1保存的epoch0/1共40项完整输出均完成比较，均未通过既定数值门槛。源码提交、冻结脚本、输入文件哈希和计时模式一致；GPU分配分别为0–3与4–7，不能把差异归因于某一后端，也不能单独归因为进程随机性。第二次计时运行的完整样本审核独立进行。

|保存阶段|输出|指标|最小观测值|最大观测值|通过数|
|---|---|---|---:|---:|---:|
|warmup|latent|relative_l2|0.024312|0.206125|0/10|
|warmup|pixels|rmse|0.015928|0.141053|0/10|
|steady|latent|relative_l2|0.024312|0.206125|0/10|
|steady|pixels|rmse|0.015928|0.141053|0/10|

像素先归一化到[-1,1]再计算RMSE；latent使用相对L2及原allclose门槛。没有临时放宽阈值。这是同源重复运行的数值差异，性能重复不构成数值等价证明；不同GPU组也是后续性能比较需保留的条件。证据：analysis/tail_480_r0_r1_saved_output_comparison.json与tail_480_r0_r1_output_summary.json。

## 六轮扩样：进程内输出哈希一致性

六次独立运行分别检查epoch 1–250的10个chunk：latent与解码输出共120组，每组250条，测量epoch完整且每组只有一个哈希。源码中哈希来自对应CPU张量的完整字节；预热epoch 0不计入该检查。

这支持每次启动内部对应chunk的记录输出重复一致；不同启动保存张量仍未通过原数值门槛。不能把进程内一致性当成跨启动或后端间等价，也不能据此确定随机数、算法选择或启动状态中的具体原因。未保存的完整输出无法从哈希还原或计算逐像素误差。

证据：analysis/tail_six_launch_hash_repeatability.json，保留每个chunk的哈希、计数及六份原始summary的SHA256；复核脚本scripts/audit_tail_hash_repeatability.py。

## 跨启动差异：随机状态与条件张量诊断

两次独立诊断采用480×832、原始参考代码、seed42、四卡Ulysses4，每次3个session。只保存各rank前12次block调用的条件、相机、文本嵌入、显式生成器状态、全局CUDA随机状态、四份噪声和输出；同步及保存会改变耗时，不计入干净性能基线。

|运行|最近保存的manifest状态|记录文件数|
|---|---|---:|
|rng_diagnostic_r0_v2|timing_complete|4|
|rng_diagnostic_r1_v2|timing_complete|4|

manifest状态不是进程存活证明。第一版钩子因继承的分辨率条件未启用，退出0但没有随机状态数据，已明确判为采集失败并保留原始目录。第二版修正入口，结束后要求四个rank都有记录；完整张量校验与两次比较结果见诊断结果节；若该节尚未生成，则比较仍待完成。

当前不能确定跨启动差异原因；全局随机状态不同本身也不能证明显式生成器产生的噪声不同。证据：analysis/rng_diagnostic_protocol.json、analysis/rng_diagnostic_r0_v1_capture_failure.json。

## 跨启动诊断结果：已记录输入一致，block输出不同

两次诊断均返回0，使用相同源码提交、冻结脚本、输入哈希和物理GPU4–7。每次四个rank各12个block、11类张量，共528个保存张量，逐一通过形状、dtype、有限值与内容哈希检查。以下逐位比较不使用放宽阈值。

|记录对象|对应张量数|逐位一致数|
|---|---:|---:|
|camera|48|48|
|condition|48|48|
|generator_after|48|48|
|generator_before|48|48|
|global_cuda_rng_before|48|0|
|noise_00|48|48|
|noise_01|48|48|
|noise_02|48|48|
|noise_03|48|48|
|output|48|0|
|prompt_embeds|48|48|

48组block输出相对L2范围0.016396–0.189984；首个block的四个rank即出现差异。相机、图像条件、文本嵌入、显式生成器前后状态及四份噪声全部相同。全局CUDA随机状态不同，但已采噪声相同，因此不能归因为输入噪声不同。

差异位于已记录的block输入到输出之间。模型权重、缓存、调度参数和逐步中间激活没有全部逐项比较；不能仅凭此归因于某个kernel、编译器或随机状态。该诊断仅覆盖480×832的两次启动及每rank前12个block，不追溯证明历史运行的输入相同，也不是优化候选验收。

原始张量：raw/rng_diagnostic_r0_v2/rng_capture与raw/rng_diagnostic_r1_v2/rng_capture。完整逐张量误差、路径和哈希：analysis/rng_diagnostic_v2_comparison.json。

## 细化诊断：加载权重一致，首次Transformer输出不同

两次480×832诊断均完成，各保存1728个张量（四rank×12 block×36类），所有文件通过内容哈希、形状、dtype与有限值检查。每个rank的1267项Transformer参数及buffer共5068对，跨启动形状、dtype、SHA256全部一致。

|首次Transformer调用的记录对象|对应张量数|逐位一致数|
|---|---:|---:|
|transformer_00_input_hidden_states|48|48|
|transformer_00_input_timestep|48|48|
|transformer_00_input_encoder_hidden_states|48|48|
|transformer_00_input_camera_hidden_states|48|48|
|transformer_00_output|48|0|

首次Transformer输出相对L2范围0.006816–0.145771。首个block的四个rank也已不同；随后Transformer调用的hidden_states不同，符合差异沿后续步骤传播的观察。

差异范围缩小到首次Transformer调用：已记录权重和四类张量输入一致，输出不同。cache对象、非张量参数及全部内部执行状态未完整比较，不能据此证明具体kernel、编译器或缓存是根因。前向钩子和CPU保存可能改变编译与执行，该结果仅用于诊断，不用于速度结论，也不代表历史基线全部同样。退出阶段worker超时与共享资源清理警告已保留。

证据：analysis/rng_diagnostic_v3_comparison.json、analysis/rng_diagnostic_v3_cross_launch_weights.json；原始张量在raw/rng_diagnostic_r0_v3与raw/rng_diagnostic_r1_v3。

## 逐层诊断v4：第0层输出首次出现差异

两次480×832四卡诊断均退出0，各2732份张量记录通过内容哈希、形状、dtype及有限值校验。加载参数和buffer共5068对哈希一致。新增层边界仅覆盖首个已捕获Transformer调用，不代表全部调用。

|第0层边界|对应数|逐位一致数|
|---|---:|---:|
|layer_block00_in_arg0|4|4|
|layer_block00_in_arg2|4|4|
|layer_block00_in_arg3|4|4|
|layer_block00_in_rotary_emb_0|4|4|
|layer_block00_in_rotary_emb_1|4|4|
|layer_block00_out_0|4|0|

进入第0层前的patch及相机投影模块输入、输出全部逐位一致。arg0为隐藏状态，arg2为时间步投影，arg3为相机状态；两个rotary记录是旋转位置编码。第0层输出在四rank均不同。

|rank|第0层输出相对L2|最大绝对误差|
|---|---:|---:|
|0|0.000325848|0.375000|
|1|0.000104068|0.015625|
|2|0.000296454|0.312500|
|3|0.000072479|0.015625|

差异定位到第0层内部或其读取的状态，尚未定位到具体kernel。缓存对象及非张量参数未记录；文本K/V在进入本次捕获调用前生成，所以没有text_embedding钩子事件。不能把已记录输入相同当作全部执行状态相同。退出清理超时警告已保留，钩子和CPU保存会改变执行，该诊断不作为性能基线或优化验收。

证据：analysis/rng_diagnostic_v4_comparison.json、analysis/rng_diagnostic_v4_cross_launch_weights.json、analysis/rng_diagnostic_v4_cache_source_review.json；原始数据：raw/rng_diagnostic_r0_v4与raw/rng_diagnostic_r1_v4。

## 第0层内部诊断v5：self-attention输出出现差异

两轮各2852份张量均通过内容哈希、形状、dtype及有限值检查；5068对参数及buffer哈希一致。首个已捕获Transformer调用的第0层中，norm1输入输出一致，self-attention已记录输入一致而输出不同。

|记录对象|对应数|逐位一致数|
|---|---:|---:|
|layer_block00sub_norm1_out|4|4|
|layer_block00sub_self_attn_in_arg0|4|4|
|layer_block00sub_self_attn_in_rotary_emb_0|4|4|
|layer_block00sub_self_attn_in_rotary_emb_1|4|4|
|layer_block00sub_self_attn_out|4|0|
|cache_block00_cross_cache_key|4|4|
|cache_block00_cross_cache_value|4|4|
|cache_block00_self_cache_block_table|4|4|
|cache_block00_self_cache_video_slots|4|4|
|cache_block00_self_cache_seq_lens|4|4|

其余已记录缓存张量与逻辑标量参数在四rank也一致；相机分支各模块输入输出一致。self-attention之后的norm3、cross-attention及FFN输入已经不同，不能把它们的输出差异当作独立根因。

|rank|self-attention输出相对L2|最大绝对误差|
|---|---:|---:|
|0|0.000318143|0.125000|
|1|0.000254624|0.125000|
|2|0.000272049|0.031250|
|3|0.000351501|0.031250|

该模块包含QKV投影、归一化、旋转位置编码、通信、分页Attention及输出投影。尚未逐一比较这些内部边界，完整key_pool/value_pool也未保存，因此不能直接归因于Attention kernel。编译跟踪、重新编译及退出清理警告已保存；钩子和CPU写出会改变执行，该结果只用于诊断，不作为性能基线。

证据：analysis/rng_diagnostic_v5_comparison.json、analysis/rng_diagnostic_v5_cross_launch_weights.json、analysis/rng_diagnostic_v5_cross_launch_cache_metadata.json。原始数据：raw/rng_diagnostic_r0_v5与raw/rng_diagnostic_r1_v5。

## 内部边界诊断v6：插桩后差异移至第1层

两轮各2908份张量通过内容哈希、形状、dtype及有限值校验。此次新增第0层self-attention内部钩子后，该层输出和分页接口输出逐位一致，但四rank的首次层输出差异均出现在第1层。最终48组block输出仍全部不同；没有解决跨启动差异。

|记录对象|对应数|逐位一致数|
|---|---:|---:|
|layer_selfdetail_qkv_out_0|4|4|
|layer_selfdetail_norm_q_out|4|4|
|layer_selfdetail_norm_k_out|4|4|
|paged_first_query|4|4|
|paged_first_key|4|4|
|paged_first_value|4|4|
|paged_first_output|4|4|
|paged_first_visible_key_after|4|4|
|paged_first_visible_value_after|4|4|
|layer_selfdetail_o_out|4|4|
|layer_block00_out_0|4|4|
|layer_block01_in_arg0|4|4|
|layer_block01_out_0|4|0|
|output|48|0|

该观察表明差异定位对插桩敏感。新增钩子改变编译边界、执行和同步，不能把v5的self-attention输出差异直接归因于其Attention kernel，也不能把v6的局部一致解释成修复成功。

分页接口返回后的有效K/V相同，不证明kernel读取瞬间的内存状态。5068对记录权重哈希与四rank逻辑缓存标量参数一致。后续应采用预先固定的编译/执行对照，验证未插桩跨启动重复性；本节不给出性能加速结论。

证据：analysis/rng_diagnostic_v6_comparison.json、analysis/rng_diagnostic_v6_cross_launch_weights.json、analysis/rng_diagnostic_v6_cross_launch_cache_metadata.json。原始数据：raw/rng_diagnostic_r0_v6与raw/rng_diagnostic_r1_v6。

## 未加诊断钩子的执行模式对照

最近记录状态：four_runs_completed; eager_40_of_40_exact; compiled_0_of_40_pass。固定480×832、TP1/SP4、seed42，每次两轮，按eager、compiled、compiled、eager顺序独立启动；只通过gpu run顺序使用四卡。

eager同时关闭模型编译和CUDA Graph预热路径，所以这是组合模式对照，不能单独归因于编译器。分别比较同模式跨启动和不同模式输出；属于小样本重复性检查，不宣称大样本性能结论。证据：analysis/execution_control_protocol_v1.json；raw/execution_control_*_v1。

## 执行模式首组比较：跨模式数值检查未通过

eager r0与compiled r0均结束，源码、输入哈希、冻结脚本、overlay及物理GPU分配一致。40对完整latent/uint8输出均完成比较，通过原定阈值的数量为0/40。

这只是不同执行模式之间的差异；两种模式各自的跨启动重复性见同模式对照节，不能仅凭本跨模式结果判断哪种模式正确或稳定。暂不把两种输出当作数值等价，也不提出可接受的加速结论。证据：analysis/execution_control_eager0_compiled0_comparison.json与analysis/execution_control_first_pair_provenance.json。

## Compiled模式跨启动重复检查

两次未加诊断钩子的compiled运行均结束，40对完整输出全部完成比较，逐位一致0/40，原定数值阈值通过0/40。源码、输入哈希、脚本、overlay、物理GPU及执行模式已核对一致。

这复现了该compiled配置的跨启动数值差异，仍不证明编译器是单独根因；eager跨启动结果见四次对照汇总。证据：analysis/execution_control_compiled_repeat_comparison.json、analysis/execution_control_compiled_pair_provenance.json。

## 四次执行模式对照完成：eager可重复，compiled仍有差异

按eager、compiled、compiled、eager顺序完成四次独立启动，每次两轮，固定480×832、TP1/SP4、seed42和GPU4–7；没有诊断张量钩子或profiler。每次40份完整latent/uint8输出保留。

|同模式跨启动比较|逐位一致|通过原定数值阈值|
|---|---:|---:|
|eager|40/40|40/40|
|compiled|0/40|0/40|

eager的40对输出最大绝对误差为0。该小样本控制将差异范围缩小到编译/CUDA Graph组合执行路径相关行为；enforce_eager同时改变两项，不能单独归因于编译器或CUDA Graph，也不能证明其他形状、模型或更多启动都稳定。跨模式输出已未通过原定阈值，不能把eager当作与compiled数值等价的替代，更不能直接据此宣称可接受加速。

证据：analysis/execution_control_eager_repeat_comparison.json、analysis/execution_control_compiled_repeat_comparison.json及两个pair_provenance.json；运行与全部输出：raw/execution_control_*_v1。

## 保留编译、关闭CUDA Graph配置的补充对照

最近记录状态：both_completed; default_repeat_0_of_40_pass; disabling_cudagraph_configuration_not_sufficient。采用compile mode=default，并显式设置TORCHINDUCTOR_CUDAGRAPHS=0，仍编译模型且不加诊断钩子；计划两次独立启动，每次两轮。

本机PyTorch模式定义中，reduce-overhead相对default增加triton.cudagraphs=True。已有compiled日志确认模型没有rollout预热请求，因此此前不能把该预热视为已实际执行。该补充对照检验CUDA Graph配置与跨启动差异的关系；配置不等于逐kernel运行时追踪证明。源码快照：raw/compile_mode_source_review_v1；协议：analysis/execution_control_default_protocol_v1.json。

## Eager与default编译首组输出比较

40对完整输出全部完成比较，原定数值阈值通过0/40。源码、输入、冻结脚本和GPU分配已核对一致。

关闭CUDA Graph配置并不自动使compiled输出与eager数值等价。default模式自身的跨启动重复性见补充对照结果；本跨模式结果不能替代该检查。证据：analysis/execution_control_eager0_default0_comparison.json、analysis/execution_control_eager_default_pair_provenance.json。

## 关闭CUDA Graph配置后仍有跨启动差异

两次default编译运行均结束，40对完整输出全部比较，逐位一致0/40，原定阈值通过0/40。四进程日志确认mode=default，运行环境显式TORCHINDUCTOR_CUDAGRAPHS=0。

这说明关闭CUDA Graph配置不足以消除该案例的跨启动差异，不能将问题仅归因于reduce-overhead的CUDA Graph路径。当前小样本中eager可逐位复现，两个编译模式均未通过跨启动阈值；具体编译变换、算子实现或状态交互仍未定位，也未用本次运行时trace证明所有CUDA Graph调用缺席。

证据：analysis/execution_control_default_repeat_comparison.json、analysis/execution_control_default_pair_provenance.json；完整运行：raw/execution_control_default_r0_v1与raw/execution_control_default_r1_v1。

## Eager参考下的cuDNN cross-attention整模型验证

最近记录状态：completed; 40_of_40_outputs_failed_original_threshold; candidate_not_accepted。已验证eager参考两次启动40对完整输出逐位一致，现在只更换cross-attention为CUDNN_ATTN，保持self-attention、480×832、TP1/SP4、seed42和两轮协议。

先核对后端实际分派，再比较完整输出。沿用原定阈值，不因算子回放通过而宣称整模型等价；数值检查通过前不接受速度优化结论。证据：analysis/cross_candidate_eager_protocol_v1.json；raw/cross_candidate_eager_r0_v1。

## Eager整模型cuDNN cross候选未通过数值验收

候选运行退出0，四rank日志选择CUDNN_ATTN，脚本仅改变cross后端配置。与已验证可重复的eager参考比较，40对完整输出逐位一致0/40、原定阈值通过0/40。
latent相对L2范围为0.034454–0.261386。

候选不能按当前标准接受。算子回放的局部误差通过没有转化为整模型数值等价，不能将本候选宣称为有效加速。后端选择日志不代替本次逐kernel trace；原始输出、计时和日志全部保留以便继续分析。证据：analysis/cross_candidate_eager_r0_comparison.json、analysis/cross_candidate_eager_r0_provenance.json；raw/cross_candidate_eager_r0_v1。

## 执行模式与候选：小样本计时记录

每次仅取epoch1的CPU-ready交付间隔5–8，共4个间隔；排除编码和网络。没有profiler或诊断钩子，但仍有正常基准记录和输出保存。这里只报告实测，不计算小样本p99或宣称稳定加速。

|运行|交付FPS|平均交付间隔ms|平均DiT ms|
|---|---:|---:|---:|
|execution_control_eager_r0_v1|11.687|1026.822|861.749|
|execution_control_compiled_r0_v1|14.256|841.734|679.137|
|execution_control_compiled_r1_v1|14.334|837.179|672.132|
|execution_control_eager_r1_v1|11.700|1025.628|859.806|
|execution_control_default_r0_v1|14.394|833.679|659.429|
|execution_control_default_r1_v1|14.516|826.660|661.313|
|cross_candidate_eager_r0_v1|11.750|1021.269|860.047|

跨模式/候选数值验收未通过，所以表中较快结果不能作为可接受优化收益。DiT时间和交付间隔具有不同边界，不能相加。完整逐间隔数据、源summary路径和哈希：analysis/execution_control_pilot_timings.json。

## 扩样输出的保存范围

当前冻结计时协议保存预热epoch=0和首个测量epoch=1的完整输出：共20个latent及20个uint8像素张量，文件已核对存在。后续epoch保留逐块latent哈希与计时日志，成功结束的summary还包含解码哈希与计时，但没有保存每轮完整latent/像素张量。

因此后续可重算已记录计时、检查哈希是否一致，并比较前两轮完整输出；不能从哈希恢复其他轮次图像或补算其逐像素误差。若需要所有轮次完整输出，需要额外采集。当前运行没有中途加入磁盘写入改变计时协议。证据：analysis/tail_output_persistence_scope.json。

## 历史记录：定时写出v1尾部仍不完整

NVTX时钟锚点与原映射偏差约-67.84微秒，锚点偏移范围约59.85微秒。最后一块仅一个DiT进程有具名Attention记录，其余三个没有；因此本次定时写出未解决覆盖问题，不采用其稳态kernel统计作为完整归因。原始报告及流文件均已保存。后续worker显式profiler stop的v2已完成并通过时钟与尾部存在性检查，详见v2专节；本节v1仍不用于完整稳态归因。证据：analysis/original_nsys_flush_v1_anchor_audit.json；后续任务：raw/worker_flush_trace_workflow_v2.json。

## Worker显式写出补采：尾部覆盖审核结果

状态：clock_and_tail_checks_passed。独立NVTX锚点与会话映射偏差-34.75微秒，锚点偏移范围77.28微秒。各worker在末块同步后显式profiler stop，本次最后一块记录如下。

|DiT进程|最后块kernel次数|其中具名Attention次数|
|---|---:|---:|
|341867|8692|400|
|341868|8692|400|
|341869|8692|400|
|341870|8692|400|

四进程最后块均有记录，解决本次补采的尾部缺失检查；这不追溯修复旧报告，也不证明全部层/shape的每次调用完整。退出仍有worker超时警告，不能仅凭退出状态推断采集质量。额外同步及写出使本次属于诊断采集，不能作为未插桩性能或与旧时间线直接比较加速。物理GPU明细在analysis/nsys_worker_flush_v2_physical生成；证据：analysis/original_nsys_worker_flush_v2_anchor_audit.json、raw/original_nsys_worker_flush_v2/pipeline.nsys-rep。

## 新时间线：逐块覆盖与物理GPU活动

两轮共20块×4个DiT进程，80个进程/块组合均有400次具名Attention记录，且每块映射到四张不同物理GPU。此项检查通过：True。它验证逐块记录存在与映射，尚非逐层shape完整性证明。完整记录共1091749次kernel，包含初始化与预热。

以下窗口为epoch=1的CPU-ready交付间隔5–8，合计约3.411秒；属于插桩诊断。各列为同卡区间并集或交集，不能跨列或跨卡直接相加为墙钟。

|物理GPU|非NCCL kernel并集ms|NCCL并集ms|两者交集ms|已记录kernel/拷贝并集ms|未记录活动ms|
|---|---:|---:|---:|---:|---:|
|0|2685.68|339.03|82.00|2995.72|414.83|
|1|2684.93|328.39|82.18|2976.85|433.70|
|2|2698.23|322.07|79.49|2986.39|424.16|
|3|2720.95|277.65|61.82|2982.63|427.92|

非NCCL kernel包含计算、转换等，不能全部称为计算；交集不证明SM资源同时执行或净收益。未记录活动不能直接归因为CPU瓶颈。

|拷贝方向|调用数|累计GPU ms|相交调用完整字节数|
|---|---:|---:|---:|
|Device-to-Host|44|7.486|62300196|
|Host-to-Device|3632|6.332|22254464|
|Device-to-Device|37008|176.302|295678976000|

字节数属于与窗口相交的完整调用，不按边界截断；累计拷贝时间不是墙钟。原始Nsight、SQLite及逐kernel gzip明细均保留。证据：analysis/worker_flush_v2_block_coverage.json、analysis/nsys_worker_flush_v2_physical/summary.json与kernel_invocations.jsonl.gz。

## 新时间线的提交调用关联

按完整进程编码前缀加correlationId关联，共1091749条已记录kernel，各有唯一提交API匹配，无歧义或缺失匹配。这个分母只包含已记录事件，不表示真实执行的所有kernel必定被采集。

|提交API|关联kernel数|
|---|---:|
|cudaLaunchKernel_v7000|371513|
|cuLaunchKernelEx|19776|
|cudaLaunchKernelExC_v11060|18184|
|cuLaunchKernel|16996|
|cudaGraphLaunch_v10000|665280|

同一次CUDA Graph提交可对应很多kernel；表中是kernel数量，不能当作API调用次数。API持续时间不是GPU执行时间、排队时间或关键路径阻塞时间。该关联提供后续逐调用归因入口，不包含融合kernel的完整shape。证据：analysis/worker_flush_v2_runtime_link_summary.json与analysis/nsys_worker_flush_v2_physical/kernel_runtime_links.jsonl.gz。

## 新时间线的主要累计耗时类别

按完整名称规则分类，窗口仍为epoch=1交付间隔5–8。分母是四卡累计kernel时间，包含跨卡及同卡重叠；不是3.411秒墙钟，也不是硬件利用率。名称分类不能替代源码层及shape关联。

|类别|名称数|调用数|累计GPU ms|占累计kernel时间|
|---|---:|---:|---:|---:|
|矩阵乘具名kernel|11|32848|4358.40|36.12%|
|Attention具名kernel|3|6448|3335.48|27.64%|
|NCCL通信|3|14592|1267.15|10.50%|
|索引与拷贝具名kernel|16|34808|896.90|7.43%|
|cuDNN卷积相关kernel|8|1776|850.86|7.05%|
|Triton融合kernel|14|51200|675.76|5.60%|
|其他名称|44|30320|525.98|4.36%|
|布局转换具名kernel|3|4992|155.49|1.29%|

按这一诊断口径，矩阵乘与Attention是最大的两个累计耗时类别，可优先补齐对应shape及NCU证据；尚不能据此认定计算受限或给出端到端加速上限。索引/拷贝kernel与CUPTI memcpy记录不同；NCCL可能与其他活动重叠，不能把其累计时长全部当作可消除等待。完整名称成员及规则见analysis/worker_flush_v2_kernel_families.json与scripts/summarize_kernel_families.py。[下载分类CSV](/worker_flush_v2_kernel_families.csv)。

## VAE精度口径：FP32张量不等于严格FP32计算

本报告的FP32 VAE指模型/输入张量精度。新诊断稳态窗口包含1680次、7种名称明确含TF32的cuDNN卷积kernel；独立真实VAE算子回放记录cudnn_allow_tf32=true，同时matmul_allow_tf32=false。这两项策略不能混为一谈。

因此当前VAE不能描述为“全部使用严格IEEE FP32计算”。也不能推断每个VAE操作都使用TF32；其他算子仍有不同路径。报告没有做禁用cuDNN TF32的整模型对照，不能量化其速度或输出差异。计算上限及后续NCU分析必须匹配实际算术路径，不能套用FP32 IEEE GEMM校准值。证据：analysis/vae_precision_evidence.json，含完整kernel名称。

## VAE逐块记录覆盖

20次decode的start到CPU-ready窗口×4个VAE进程，共80个组合，均记录到具名cuDNN卷积相关kernel（包含工作区初始化），并映射到每轮四张不同物理GPU。8个组合记录114次，其余72个记录111次；这些是名称筛选后的调用数，不等于全部VAE算子数。该项验证逐块卷积记录存在，不能证明全部层shape完整或数值正确。证据：analysis/worker_flush_v2_decode_coverage.json。

## 历史记录：高分辨率纯DiT v1保存额度不足

704×1280 TP1/SP1纯DiT完成十块生成，但代表输入仅保存3/12个。补采协调脚本漏传原矩阵使用的12 GiB额度，触发采集器默认2 GiB上限；这次未进行回放，不能增加审核点数。后续v2已显式设置12 GiB，在独立目录保存并审核通过12/12代表输入；旧记录保留。实测元数据表明六个self代表输入合计约9.52 GB，另需保存cross输入。证据：analysis/dit_capture_budget_failure_v1.json；后续状态：raw/dit_recovery_workflow_v2.json。纯DiT成功不等于此前含VAE的流水线失败已解决。

## 修复后的算子输入采集

v2已经保存66条实际算子输入记录，工作负载状态：completed。这是输入证据，不能当作硬件分析完成。

|阶段|实际ATen调用|保存记录数|
|---|---|---:|
|dit|aten.linear.default|32|
|vae|aten.conv3d.default|32|
|output|aten.to.dtype|2|

完整性审核：通过。

普通用户单卡回放状态见raw/operator_replay_workflow_v2.json；NCU计数器采集仍受调度身份问题限制。证据：analysis/operator_capture_v2_inventory.json。

## 真实算子单卡回放结果

证据审核状态：evidence_checks_passed。共18项rank0代表输入，输入哈希、有限值和实际GPU kernel均已核对。非逐位一致输出数：0。这些是原ATen调用的回放，不是新实现验收，也没有NCU硬件计数器。

|阶段|实际调用|输出形状|逐位一致|记录kernel数|
|---|---|---|---|---:|
|dit|aten.linear.default|[4680, 5120]|True|1|
|dit|aten.linear.default|[1, 4680, 5120]|True|1|
|dit|aten.linear.default|[1, 4680, 5120]|True|1|
|dit|aten.linear.default|[3, 5120]|True|1|
|dit|aten.linear.default|[3, 5120]|True|1|
|dit|aten.linear.default|[1, 3, 30720]|True|1|
|dit|aten.linear.default|[1, 1170, 15360]|True|1|
|dit|aten.linear.default|[1, 1170, 5120]|True|1|
|vae|aten.conv3d.default|[1, 16, 3, 60, 104]|True|2|
|vae|aten.conv3d.default|[1, 384, 1, 60, 26]|True|4|
|vae|aten.conv3d.default|[1, 384, 1, 60, 26]|True|5|
|vae|aten.conv3d.default|[1, 384, 1, 120, 52]|True|2|
|vae|aten.conv3d.default|[1, 384, 1, 120, 52]|True|6|
|vae|aten.conv3d.default|[1, 384, 1, 120, 52]|True|6|
|vae|aten.conv3d.default|[1, 192, 1, 240, 104]|True|5|
|vae|aten.conv3d.default|[1, 96, 1, 480, 208]|True|5|
|output|aten.to.dtype|[1, 3, 9, 480, 832]|True|1|
|output|aten.to.dtype|[1, 3, 12, 480, 832]|True|1|

[下载输入形状、完整kernel名称与证据路径 CSV](/operator_replay_v2.csv)。Profiler时长不能当作未插桩算子性能；这些输入仍需硬件计数器才能支持对应的计算/访存瓶颈归因。

## 真实算子回放与模型kernel的候选关联

18项独立算子输入中，14项至少有一个kernel在整模型记录中匹配完整名称、grid、block、每线程寄存器数和共享内存大小。8项linear全部找到匹配，8项VAE卷积中的6项找到匹配，2项uint8转换均未找到同名匹配。

这比只匹配名称提供更多线索，但同一启动配置仍可能对应多个输入shape，因此没有给整模型事件填入“确定输入shape”。未匹配也不代表整模型缺少该操作，编译融合或算法选择差异尚待进一步关联。此结果用于筛选后续归因对象，不能计算为完整shape覆盖率。全部候选与原始启动参数见analysis/operator_model_launch_candidates.json。

## 保存额度修复后的高分辨率输入

v2实际保存12/12个代表输入，完整性审核通过。旧v1的额度不足记录保留；新输入的后端回放需单独审核，不能仅凭采集成功增加已核验点数。证据：raw/capture_dit_only_704x1280_tp1_sp1_v2/capture_audit.json。

## Attention执行失败记录

以下为保存结果中的失败尝试，包含历史运行，不等于最终矩阵失败点数。运行中的结果仍会增加。OOM表示该输入与实现组合在本次显存条件下失败，不表示所有后端或整模型均不支持。失败没有有效计时，不参与最快后端比较。

|运行|后端|角色/阶段/帧|有效Q/K长度|状态|
|---|---|---|---|---|
|attention_replay_dit_only_704x1280_tp1_sp1_v2|sdpa_math|self/denoise/27|10560/63360|oom|
|attention_replay_dit_only_704x1280_tp1_sp1_v2|sdpa_math|self/commit/9|10560/42240|oom|
|attention_replay_dit_only_704x1280_tp1_sp1_v2|sdpa_math|self/denoise/9|10560/42240|oom|
|attention_replay_dit_only_704x1280_tp1_sp1_v2|sdpa_math|self/commit/27|10560/63360|oom|

完整报错、输入哈希和证据路径：analysis/attention_execution_failures.json。缺失这些组合的有效计时，会限制全后端比较；其他通过审核的组合仍可单独分析。

## 高分辨率长KV的显存边界

704×1280 TP1/SP1中，Q长度10560、40个head。SDPA math在self K=10560时完成，在K=42240和63360时，denoise与commit共4点OOM；相同输入的其他六接口全部通过独立回放审核。cross K=512的math也完成。因此这是本次长KV与math实现组合的容量失败，不能泛化为整组配置不支持。

|阶段/帧|有效K|math结果|单份FP32稠密分数矩阵GiB（计算值）|报错申请GiB（实际记录）|
|---|---:|---|---:|---:|
|commit/0|10560|measured|16.62|无失败申请|
|commit/9|42240|oom|66.47|16.62|
|commit/27|63360|oom|99.70|99.7|
|denoise/0|10560|measured|16.62|无失败申请|
|denoise/9|42240|oom|66.47|16.62|
|denoise/27|63360|oom|99.70|99.7|

矩阵大小按H×Q×K×4计算，只解释稠密表示的规模，不是测得的峰值，也未证明报错申请对应哪一个中间张量。K=42240时失败申请16.62 GiB，与单份分数矩阵66.47 GiB不同，更不能将一次申请当作总显存。容量失败点不参与七接口完整排名；其他接口成功也不证明含VAE整流水线可运行。证据：analysis/attention_capacity_boundary.json，包含完整错误与输入哈希。

## 实验身份与口径

算子采集v1没有生成矩阵乘、VAE卷积或输出转换输入，不能视为硬件分析完成。CPU诊断确认旧过滤器遗漏了推理模式下的aten.linear、aten.conv3d和aten.to.dtype复合调用，修复后的过滤器已在v2实际GPU运行中生成输入，完整性及回放状态见下方。证据：analysis/operator_capture_v1_failure.json。v2补采由raw/operator_recovery_workflow_v2.json记录；高分辨率纯DiT补采由raw/dit_recovery_workflow_v1.json记录，均串行等待、不预约GPU。

高分辨率TP1/SP1采集在2026-09-09 20:12出现显存分配失败警告。20:14现场记录DiT与一个VAE进程共享物理GPU，分别占用136936 MiB和5906 MiB；当时任务尚未退出。这是运行中显存现场，不是终态OOM判定或独立算子峰值。证据：analysis/highres_tp1_sp1_memory_observation.json及其原始日志；该次任务随后于20:17失败退出（返回码1）：VAE NCCL超时后子进程SIGABRT，未完成完整输入采集及回放。终态证据：analysis/highres_tp1_sp1_failure.json。不能据此宣称所有TP1/SP1 Attention运行都不支持。

原始源码 b9cdea7b6da42d51cd5772548c616655fec21668 加冻结补丁；rebase c59234d1974fe280116c1b77473d110ab53766e3。基线为 480×832、BF16 DiT、FP32 untiled Wan VAE、4 DMD steps 与 clean KV commit。四卡顺序执行 Ulysses4 DiT 和宽度空间分片 VAE；不是权重 TP4 解码。
计时边界为 rank0 CPU-ready uint8，包含相关传输与转换，不包含后续 RGB IPC、编码、网络或客户端播放。已有六次 clean baseline 每次仅80个稳态间隔，不足最终尾延迟验收。

## 已有基线与时间线

补充同版本控制：original_r0与original_r1的20个已保存latent/像素文件均完成比较，均未通过候选协议的严格数值门槛。逐chunk latent相对L2为3.33%–17.45%；归一化像素RMSE为0.02166–0.10453。同文件身份对照20/20通过。证据：analysis/original_repeat_candidate_gate_control.json。此处是同版本基线控制，不是候选比较。基线跨启动差异的来源仍需隔离，不能将差异全部归因于换后端，也不能因此临时放宽门槛。

补充epoch=1的保存输出检查：original_r0与original_r1的10个steady latent及10个steady pixel全部完成比较，全部未通过原门槛；latent relative L2为3.745%–45.339%，归一化像素RMSE为0.02454–0.33010。这里的steady文件保存整个epoch=1的10个chunk，不仅是计时使用的交付间隔5–8，也不是全部20个测量session。原始比较见analysis/original_repeat_candidate_gate_control_with_steady.json。候选比较已扩展到预热及epoch=1共40个文件。

基线身份复核：original_r0/r1的模型输入哈希、代码及冻结脚本哈希、seed=42、分辨率和GPU4–7分配一致；记录中的环境差异为各次运行的overlay/cache目录。尚未保存可用于逐步对照的运行时RNG状态和全部中间张量，不能据此定位数值差异来源。证据：analysis/original_repeat_provenance_control.json。

进程内重复性检查：六次基线分别核对20个测量session×10个chunk，全部在同一运行、相同chunk上得到唯一latent哈希；保存的epoch=1张量原始字节哈希也全部匹配日志。这支持已记录latent在进程内重复稳定，跨启动差异应优先检查启动相关状态或算法选择，但尚未证明具体原因。此结论依据逐session日志及epoch=1张量，不表示保存了每个session的完整张量，也不构成像素逐位一致证明。证据：analysis/session_repeatability.json。

原版三次 steady FPS 为14.425、14.003、14.392；rebase为14.405、14.430、14.351。完整输出等价性未通过验收，同版本跨启动也观察到数值差异，不能将所有差异归因于rebase。

Nsight Systems 原版保存954123次 kernel 调用，新版保存988240次；均为完整进程记录，包含启动、预热、验证和清理，不能直接与稳态计时比较。analysis/nsys_*_physical 提供完整调用明细及完整/测量/稳态窗口分组。按(pid,cudaId)映射物理GPU，旧的未映射重叠分析已排除。
kernel 名称分组尚未全部关联实际shape、层、阶段；因此完整shape组覆盖率和逐组MFU仍未完成。累计线程等待和多卡GPU时长不是墙钟时间；时间线交集也不能证明硬件资源并行或收益。

## 独立Attention的局部最快结果

仅纳入七种接口各有三组重复且已审核的相同输入，共140个上下文。按三组mean的中位数选最低者；重复范围不重叠仅是描述性检查，不是统计显著性或独立启动置信区间。

|角色|最低中位数接口|上下文数|与全部竞争者重复范围不重叠|
|---|---|---:|---:|
|self|native_dense|55|29|
|self|vllm_fa3|13|2|
|cross|native_dense|6|6|
|cross|sdpa_cudnn|63|57|
|cross|sdpa_flash|2|0|
|cross|vllm_fa2|1|1|

Native dense与varlen FA3为两个接口，不能当作不同硬件架构。cross的cuDNN局部胜出不改变整模型数值门槛未通过的事实，不能直接推荐替换。

![局部最快接口与缺失配置](/attention_winners.png)

图中星号表示与至少一个竞争者的重复范围重叠；灰格包括缺失或因OOM未满足七接口比较条件的上下文；OOM明确标注。每个单元格固定本配置的真实输入，不比较不同TP/SP行的端到端收益。[下载逐点结果 CSV](/attention_winners.csv)。完整证据：analysis/attention_winners.json。

## 候选实际 kernel 证据

按完整时间线统计，包含初始化、预热与清理；调用数和累计多卡时长不能直接解释为稳态墙钟收益。以下仅统计名称明确含native_sdpa的cuDNN Attention kernel，未将VAE的cuDNN卷积计入。

|运行|全部kernel调用|cuDNN native SDPA调用|
|---|---:|---:|
|integrated_cross_diagnostic_default_v1|954202|0|
|integrated_cross_diagnostic_cudnn_attn_v1|1034741|13642|

证据：analysis/integrated_dispatch_audit_v1.json保存完整kernel名称、计数与进程/GPU映射。两次完整调用数不同，不能将总耗时直接相除作为后端加速比。

## 整模型 cross-attention 候选诊断

默认后端与cross=CUDNN_ATTN两次Nsight Systems诊断均已结束。保存输出共40个比较，通过0个；候选未通过预设整模型数值门槛。两次均开启profiler，不能作为未插桩性能A/B结果。基线跨启动数值差异尚未定位，不能把全部误差归因于后端。实际时间线已核验出现cuDNN SDPA kernel；逐调用与层/角色的完整关联仍未完成。

|阶段|输出|通过/总数|relative L2范围|RMSE范围|
|---|---|---:|---:|---:|
|warmup|latent|0/10|0.016349–0.172994|0.010632–0.147762|
|warmup|pixels|0/10|0.018224–0.182353|0.010674–0.115392|
|steady|latent|0/10|0.016511–0.182018|0.010736–0.153048|
|steady|pixels|0/10|0.016715–0.203531|0.009789–0.126815|

证据：analysis/integrated_cross_diagnostic_v1_output_comparison.json；原始输出与时间线：raw/integrated_cross_diagnostic_default_v1和raw/integrated_cross_diagnostic_cudnn_attn_v1。像素误差按[-1,1]归一化计算。

## 已采集的 Attention 元数据范围

以下为已通过输入采集审核的独立eager运行。API调用数与GPU kernel调用数不同；同一API可能产生多个kernel。完整张量仅保存代表输入，metadata覆盖的层数不能当作完整张量覆盖。

|采集配置|self API调用|cross API调用|shape/context分组|层角色前缀数|
|---|---:|---:|---:|---:|
|capture_cross_480x832_tp1_sp4_v1|16000|16000|16000|80|
|capture_dit_only_704x1280_tp1_sp1_v2|2000|2000|4000|80|
|capture_evening_480x832_tp1_sp1_v1|4000|4000|4000|80|
|capture_evening_480x832_tp2_sp1_v1|8000|8000|8000|80|
|capture_evening_480x832_tp2_sp2_v1|16000|16000|16000|80|
|capture_evening_480x832_tp4_sp1_v1|16000|16000|16000|80|
|capture_evening_704x1280_tp1_sp2_v1|8000|8000|8000|80|
|capture_evening_704x1280_tp2_sp1_v1|8000|8000|8000|80|
|capture_evening_704x1280_tp2_sp2_v1|16000|16000|16000|80|
|capture_evening_704x1280_tp4_sp1_v1|16000|16000|16000|80|
|capture_retry_480x832_tp1_sp2_v1|8000|8000|8000|80|
|capture_retry_704x1280_tp1_sp4_v1|16000|16000|16000|80|

80个层角色前缀对应40层的self/cross角色。分组包含rank和阶段上下文，不能将不同配置的分组数直接解读为覆盖率高低。详细索引：analysis/shape_inventory.json。尚未与旧Nsight Systems逐kernel事件建立可靠关联。

## 未插桩交付延迟分布

由原始CPU-ready时间戳重新计算，核对六次运行共480个间隔，均与原统计一致。逐运行展示，不把不同版本或进程混合成一次样本。每次仅80个间隔，因此P99接近样本最大值，只供描述，不能作为最终尾延迟验收。

|运行|样本数|平均ms|P50 ms|P95 ms|P99 ms|最大ms|稳态FPS|
|---|---:|---:|---:|---:|---:|---:|---:|
|original_r0|80|831.87|832.05|839.13|841.33|843.90|14.425|
|original_r1|80|856.95|858.78|880.30|889.31|902.75|14.003|
|original_r2|80|833.80|831.98|848.15|859.13|871.43|14.392|
|rebased_r0|80|833.06|833.40|840.16|844.64|845.26|14.405|
|rebased_r1|80|831.58|830.98|839.05|843.48|847.07|14.430|
|rebased_r2|80|836.16|834.93|849.81|858.57|863.86|14.351|

统计窗口为各测量session的交付间隔5–8，排除预热epoch0。每个间隔交付12帧，FPS=12/平均间隔；不包含编码、网络或客户端播放。分位数采用排序后的线性插值。

[下载480个原始时间戳差值 CSV](/clean_delivery_intervals.csv)。复核证据：analysis/clean_delivery_audit.json。

### 首块与完整rollout

每次运行20个测量session，每session117帧；保留首块fill和末块drain。这里的rollout是已记录session起止墙钟，TTFC为session开始至首块CPU-ready。两者仍不包含编码、网络或客户端播放。

|运行|TTFC P50 s|TTFC P95 s|TTFC最大 s|rollout平均 s|rollout P95 s|rollout FPS|
|---|---:|---:|---:|---:|---:|---:|
|original_r0|2.381|2.575|3.075|9.571|9.753|12.224|
|original_r1|2.415|2.554|2.582|9.817|9.995|11.919|
|original_r2|2.388|2.599|2.661|9.590|9.716|12.200|
|rebased_r0|2.382|2.427|2.653|9.545|9.611|12.258|
|rebased_r1|2.379|2.448|2.530|9.525|9.606|12.283|
|rebased_r2|2.386|2.518|2.586|9.582|9.749|12.210|

首块及rollout分布每次只有20个样本，不能替代大样本尾延迟验收。完整量化结果及与旧统计的逐字段核对见analysis/clean_delivery_audit.json。

## 实验矩阵覆盖

按12种配置、每种12个代表输入×7种后端/接口核对。这个分母只描述本轮回放矩阵，不代表全部层或硬件分析覆盖率。
已审核1004/1008个点；全部84点齐全的配置为11/12。已尝试失败与尚未尝试分开记录，不用零耗时代替失败。

|配置|已审核/预期|最近任务|任务状态|
|---|---:|---|---|
|480x832_tp1_sp4|84/84|retry-audit-480x832_tp1_sp4|completed|
|704x1280_tp1_sp4|84/84|retry-audit-704x1280_tp1_sp4|completed|
|480x832_tp1_sp2|84/84|retry-audit-480x832_tp1_sp2|completed|
|704x1280_tp1_sp2|84/84|evening-audit-704x1280_tp1_sp2|completed|
|480x832_tp2_sp2|84/84|evening-audit-480x832_tp2_sp2|completed|
|704x1280_tp2_sp2|84/84|evening-audit-704x1280_tp2_sp2|completed|
|480x832_tp2_sp1|84/84|evening-audit-480x832_tp2_sp1|completed|
|704x1280_tp2_sp1|84/84|evening-audit-704x1280_tp2_sp1|completed|
|480x832_tp4_sp1|84/84|evening-audit-480x832_tp4_sp1|completed|
|704x1280_tp4_sp1|84/84|evening-audit-704x1280_tp4_sp1|completed|
|480x832_tp1_sp1|84/84|evening-audit-480x832_tp1_sp1|completed|
|704x1280_tp1_sp1|80/84|evening-capture-704x1280_tp1_sp1|failed|

已记录尝试1008个点，其中执行失败4个；失败点不计入有效计时。

running表示任务流程尚未结束，可能包括调度等待；具体GPU进程以实时看板为准。等待超时不等于配置不支持。

## 稳态时间线的直接证据

以下均为profiler下交付间隔5–8，表中kernel与API耗时为累计值，不能相加为请求墙钟时间。不同版本的profiler结果不能替代clean A/B。

### Kernel与主机提交API的关联

按完整进程编码前缀和correlationId关联原始CUPTI runtime表，避免不同进程相同correlationId串联。所有kernel事件均有一条关联记录，不能匹配的保留missing；GPU事件本身仍已记录。

|源码|kernel总数|唯一API匹配|未匹配|关联到Graph提交的kernel数|
|---|---:|---:|---:|---:|
|original|954123|953997|126|532728|
|rebased|988240|988171|69|565194|

这里覆盖完整trace，包含初始化与预热。一个Graph提交可对应大量kernel，因此最后一列不是Graph提交次数。主机API耗时也不能直接当作GPU排队等待或关键路径。kernel/API rowid及时间保存在analysis/nsys_*_physical/kernel_runtime_links.jsonl.gz；汇总为analysis/kernel_runtime_link_summary.json。层、shape与算子源码的关联仍不完整。

### Nsys已经记录的启动资源

Nsys原始kernel事件含grid/block、每线程寄存器数和静态/动态共享内存分配。这些是启动描述，不是NCU采样的实际occupancy、访存带宽或warp stall；无需假称硬件计数器已经补齐。

|源码|启动参数分组数|稳态调用数|完整资源表|
|---|---:|---:|---|
|original|260|74756|[CSV](/launch-resources-original-steady.csv)|
|rebased|262|101350|[CSV](/launch-resources-rebased-steady.csv)|

分组键为kernel名称、grid、block、寄存器及共享内存参数，跨rank累计；并非张量shape分组。first_event_id定位对应原始SQLite的kernel rowid。调用总数已与完整稳态名称排名核对一致。

### original：耗时最多的五个名称组

|Kernel名称|调用次数|累计GPU ms|
|---|---:|---:|
|nvjet_sm90_tst_128x200_64x5_2x1_v_bz_coopA_bias_TNT|7618|917.506|
|void cutlass::device_kernel<flash::enable_sm90_or_later<flash::FlashAttnFwdSm90<flash::CollectiveMainloopFwdSm90<(int)2, cute::tuple<cute::C<(int)1>, cute::C<(int)1>, cute::C<(int)1>>, cute::tuple<cute::C<(int)128>, cute::C<(int)160>, cute::C<(int)128>>, (int)128, cutlass::bfloat16_t, float, cutlass::arch::Sm90, (bool)0, (bool)0, (bool)0, (bool)1, (bool)0, (bool)0, (bool)0, (bool)1, (bool)1, (bool)1, (bool)0, (bool)0, cutlass::bfloat16_t, (int)1>, flash::CollectiveEpilogueFwd<cute::tuple<cute::C<(int)128>, cute::C<(int)128>, cute::C<(int)160>>, cute::tuple<cute::C<(int)1>, cute::C<(int)1>, cute::C<(int)1>>, cutlass::bfloat16_t, cutlass::arch::Sm90, (int)256, (bool)1, (bool)1, (bool)0, (bool)0, (int)1>, flash::VarlenDynamicPersistentTileScheduler<(int)128, (int)160, (int)256, (int)128, (bool)0, (bool)1, (bool)1, (bool)0, (bool)0, (bool)1>>>>(T1::Params)|848|840.687|
|ncclDevKernel_SendRecv(ncclDevKernelArgsStorage<(unsigned long)4096>)|4971|454.471|
|void at::native::elementwise_kernel<(int)128, (int)2, void at::native::gpu_kernel_impl_nocast<at::native::direct_copy_kernel_cuda(at::TensorIteratorBase &)::[lambda() (instance 3)]::operator ()() const::[lambda() (instance 7)]::operator ()() const::[lambda(float) (instance 1)]>(at::TensorIteratorBase &, const T1 &)::[lambda(int) (instance 1)]>(int, T3)|13895|384.264|
|sm90_xmma_fprop_implicit_gemm_f32f32_tf32f32_f32_nhwckrsc_nhwc_tilesize256x32x32_warpgroupsize1x1x1_g1_execute_segment_k_off_kernel__5x_cudnn|288|383.824|

[下载original完整稳态名称排名](/kernel-original-steady.csv)

|拷贝方向|调用次数|累计GPU ms|相交调用完整字节数|
|---|---:|---:|---:|
|Device-to-Host|22|3.635|62300174|
|Host-to-Device|982|1.956|20017168|
|Device-to-Device|10950|53.016|90721055616|

|物理GPU|窗口ms|其他kernel ms|NCCL ms|两者交集ms|未录活动ms|
|---|---:|---:|---:|---:|---:|
|4|3408.402|1241.858|153.587|34.920|2027.217|
|5|3408.402|1033.992|83.298|15.608|2293.960|
|6|3408.402|1002.216|121.913|26.180|2297.932|
|7|3408.402|1014.352|115.955|20.021|2285.449|

|CUDA等待/同步API|次数|累计线程ms|
|---|---:|---:|
|cudaStreamSynchronize_v3020|170|1731.314|
|cudaDeviceSynchronize_v3020|33|1087.979|
|cudaStreamWaitEvent_v3020|9800|11.528|

未录活动可能涉及依赖、主机调度或未覆盖事件，尚未归因。同步API的多个线程会互相重叠；这些数字不是关键路径阻塞时间。

### rebased：耗时最多的五个名称组

|Kernel名称|调用次数|累计GPU ms|
|---|---:|---:|
|nvjet_sm90_tst_128x200_64x5_2x1_v_bz_coopA_bias_TNT|13126|1574.521|
|void cutlass::device_kernel<flash::enable_sm90_or_later<flash::FlashAttnFwdSm90<flash::CollectiveMainloopFwdSm90<(int)2, cute::tuple<cute::C<(int)1>, cute::C<(int)1>, cute::C<(int)1>>, cute::tuple<cute::C<(int)128>, cute::C<(int)160>, cute::C<(int)128>>, (int)128, cutlass::bfloat16_t, float, cutlass::arch::Sm90, (bool)0, (bool)0, (bool)0, (bool)1, (bool)0, (bool)0, (bool)0, (bool)1, (bool)1, (bool)1, (bool)0, (bool)0, cutlass::bfloat16_t, (int)1>, flash::CollectiveEpilogueFwd<cute::tuple<cute::C<(int)128>, cute::C<(int)128>, cute::C<(int)160>>, cute::tuple<cute::C<(int)1>, cute::C<(int)1>, cute::C<(int)1>>, cutlass::bfloat16_t, cutlass::arch::Sm90, (int)256, (bool)1, (bool)1, (bool)0, (bool)0, (int)1>, flash::VarlenDynamicPersistentTileScheduler<(int)128, (int)160, (int)256, (int)128, (bool)0, (bool)1, (bool)1, (bool)0, (bool)0, (bool)1>>>>(T1::Params)|1459|1428.093|
|ncclDevKernel_SendRecv(ncclDevKernelArgsStorage<(unsigned long)4096>)|7420|695.088|
|nvjet_sm90_tst_256x152_64x4_1x2_h_bz_coopA_bias_TNT|1461|388.515|
|void at::native::elementwise_kernel<(int)128, (int)2, void at::native::gpu_kernel_impl_nocast<at::native::direct_copy_kernel_cuda(at::TensorIteratorBase &)::[lambda() (instance 3)]::operator ()() const::[lambda() (instance 7)]::operator ()() const::[lambda(float) (instance 1)]>(at::TensorIteratorBase &, const T1 &)::[lambda(int) (instance 1)]>(int, T3)|13916|384.312|

[下载rebased完整稳态名称排名](/kernel-rebased-steady.csv)

|拷贝方向|调用次数|累计GPU ms|相交调用完整字节数|
|---|---:|---:|---:|
|Device-to-Host|28|4.232|62300180|
|Host-to-Device|1673|3.042|20595952|
|Device-to-Device|17742|85.043|144137728112|

|物理GPU|窗口ms|其他kernel ms|NCCL ms|两者交集ms|未录活动ms|
|---|---:|---:|---:|---:|---:|
|4|3394.550|1258.743|144.692|30.921|2000.425|
|5|3394.550|1023.110|80.731|16.239|2294.264|
|6|3394.550|2683.188|402.558|110.644|373.982|
|7|3394.550|1016.739|90.170|15.112|2290.196|

|CUDA等待/同步API|次数|累计线程ms|
|---|---:|---:|
|cudaStreamSynchronize_v3020|239|2951.602|
|cudaDeviceSynchronize_v3020|37|1488.981|
|cudaStreamWaitEvent_v3020|9832|11.522|

未录活动可能涉及依赖、主机调度或未覆盖事件，尚未归因。同步API的多个线程会互相重叠；这些数字不是关键路径阻塞时间。

## Attention：真实输入、后端、误差与计时

已核验480×832 TP1/SP4采集的元数据包含32000次Attention API调用（self/cross各16000），覆盖40层、4个rank及缓存/去噪阶段，共16000个shape/context分组。元数据覆盖与完整张量保存不同：完整Q/K/V及参考输出仅保存代表点。旧nsys逐kernel记录尚未与此独立eager运行关联，不能据此宣布kernel shape映射完成。原始索引见analysis/shape_inventory.json及其group_file。

下表仅纳入输入哈希、逐次计时和实际kernel类型审核通过的配置。每个点3组×300次CUDA event计时；重复组来自同一进程，不等同3次独立模型启动。比较保存输出及FP64抽样参考（至多32个query位置、4个head、全部有效K/V）；预声明atol=0.03、rtol=0.03、relative L2≤0.01。该检查不构成整模型等价证明。

每行固定角色、阶段和缓存位置，数值为三组mean的中位数，单位ms。不同shape行之间不直接比较后端收益。

|配置|角色/阶段/起始帧|FA2|FA3|SDPA flash|SDPA efficient|SDPA cuDNN|SDPA math|Native dense|
|---|---|---:|---:|---:|---:|---:|---:|---:|
|480x832_tp1_sp1|cross/commit/0|0.1560|0.1289|0.1672|0.3088|0.0953|2.1436|0.1219|
|480x832_tp1_sp1|cross/commit/9|0.1571|0.1295|0.1671|0.3112|0.0955|2.1448|0.1204|
|480x832_tp1_sp1|cross/commit/27|0.1572|0.1286|0.1672|0.3087|0.0947|2.1455|0.1189|
|480x832_tp1_sp1|cross/denoise/0|0.1571|0.1295|0.1672|0.3088|0.0949|2.1435|0.1213|
|480x832_tp1_sp1|cross/denoise/9|0.1573|0.1286|0.1671|0.3088|0.0948|2.1448|0.1288|
|480x832_tp1_sp1|cross/denoise/27|0.1571|0.1292|0.1671|0.3089|0.0956|2.1452|0.1231|
|480x832_tp1_sp1|self/commit/0|1.2463|0.7234|1.3310|2.6191|0.8147|17.9784|0.6860|
|480x832_tp1_sp1|self/commit/9|4.9470|2.7717|5.3070|10.3375|3.2751|70.1103|2.7732|
|480x832_tp1_sp1|self/commit/27|7.3873|4.1367|7.9554|15.3297|4.8942|105.4600|4.1254|
|480x832_tp1_sp1|self/denoise/0|1.2545|0.7337|1.3441|2.6252|0.8325|17.9769|0.7015|
|480x832_tp1_sp1|self/denoise/9|4.9365|2.7576|5.3184|10.2508|3.2926|70.1039|2.7705|
|480x832_tp1_sp1|self/denoise/27|7.3650|4.1462|7.9571|15.4722|4.9115|105.4707|4.1463|
|480x832_tp1_sp2|cross/commit/0|0.0843|0.0733|0.0892|0.1659|0.0530|1.1250|0.1158|
|480x832_tp1_sp2|cross/commit/9|0.0835|0.0727|0.0893|0.1659|0.0526|1.1251|0.1151|
|480x832_tp1_sp2|cross/commit/27|0.0833|0.0728|0.0893|0.1660|0.0528|1.1258|0.1184|
|480x832_tp1_sp2|cross/denoise/0|0.0836|0.0731|0.0893|0.1663|0.0526|1.1259|0.1250|
|480x832_tp1_sp2|cross/denoise/9|0.0834|0.0727|0.0892|0.1661|0.0538|1.1258|0.1146|
|480x832_tp1_sp2|cross/denoise/27|0.0840|0.0734|0.0895|0.1662|0.0530|1.1252|0.1162|
|480x832_tp1_sp2|self/commit/0|0.6226|0.3617|0.6705|1.3851|0.4087|9.0496|0.3382|
|480x832_tp1_sp2|self/commit/9|2.4613|1.3667|2.6639|5.3956|1.6144|35.1077|1.3414|
|480x832_tp1_sp2|self/commit/27|3.7085|2.0668|3.9722|8.1392|2.4622|52.6948|2.0382|
|480x832_tp1_sp2|self/denoise/0|0.6250|0.3667|0.6750|1.3833|0.4174|9.0499|0.3442|
|480x832_tp1_sp2|self/denoise/9|2.4692|1.3631|2.6735|5.4438|1.6346|35.1137|1.3762|
|480x832_tp1_sp2|self/denoise/27|3.7173|2.0718|3.9865|8.1435|2.4683|52.6422|2.0688|
|480x832_tp1_sp4|cross/commit/0|0.0508|0.0622|0.0494|0.0868|0.0480|0.6176|0.1103|
|480x832_tp1_sp4|cross/commit/9|0.0533|0.0665|0.0494|0.0867|0.0481|0.6178|0.1088|
|480x832_tp1_sp4|cross/commit/27|0.0505|0.0625|0.0496|0.0872|0.0480|0.6185|0.1106|
|480x832_tp1_sp4|cross/denoise/0|0.0504|0.0613|0.0504|0.0870|0.0479|0.6185|0.1145|
|480x832_tp1_sp4|cross/denoise/9|0.0508|0.0625|0.0495|0.0873|0.0478|0.6173|0.1089|
|480x832_tp1_sp4|cross/denoise/27|0.0502|0.0618|0.0496|0.0871|0.0484|0.6178|0.1111|
|480x832_tp1_sp4|self/commit/0|0.3309|0.1768|0.3567|0.6991|0.2033|4.8018|0.1679|
|480x832_tp1_sp4|self/commit/9|1.2692|0.6755|1.3729|2.7239|0.8056|18.7832|0.6673|
|480x832_tp1_sp4|self/commit/27|1.8923|1.0040|2.0524|4.1119|1.1895|28.0167|0.9872|
|480x832_tp1_sp4|self/denoise/0|0.3312|0.1859|0.3572|0.6996|0.2059|4.8040|0.1716|
|480x832_tp1_sp4|self/denoise/9|1.2679|0.6801|1.3746|2.7502|0.8077|18.7671|0.6780|
|480x832_tp1_sp4|self/denoise/27|1.8999|1.0303|2.0509|4.1231|1.2022|28.1139|0.9955|
|480x832_tp2_sp1|cross/commit/0|0.0838|0.0728|0.0892|0.1658|0.0525|1.1246|0.1159|
|480x832_tp2_sp1|cross/commit/9|0.0833|0.0726|0.0890|0.1659|0.0525|1.1244|0.1138|
|480x832_tp2_sp1|cross/commit/27|0.0839|0.0726|0.0889|0.1658|0.0527|1.1246|0.1141|
|480x832_tp2_sp1|cross/denoise/0|0.0840|0.0734|0.0891|0.1661|0.0531|1.1244|0.1199|
|480x832_tp2_sp1|cross/denoise/9|0.0839|0.0730|0.0889|0.1660|0.0525|1.1248|0.1138|
|480x832_tp2_sp1|cross/denoise/27|0.0845|0.0731|0.0897|0.1673|0.0525|1.1245|0.1178|
|480x832_tp2_sp1|self/commit/0|0.6224|0.3635|0.6710|1.3844|0.4100|9.0403|0.3375|
|480x832_tp2_sp1|self/commit/9|2.4790|1.3580|2.6619|5.3952|1.6435|35.1040|1.3688|
|480x832_tp2_sp1|self/commit/27|3.6979|2.0640|3.9728|8.0732|2.4561|52.6303|2.0574|
|480x832_tp2_sp1|self/denoise/0|0.6295|0.3688|0.6747|1.3727|0.4108|9.0404|0.3478|
|480x832_tp2_sp1|self/denoise/9|2.4556|1.3758|2.6613|5.4440|1.6203|35.1034|1.3645|
|480x832_tp2_sp1|self/denoise/27|3.7063|2.0700|3.9811|8.1448|2.4687|52.6228|2.0706|
|480x832_tp2_sp2|cross/commit/0|0.0497|0.0604|0.0491|0.0870|0.0469|0.6187|0.1082|
|480x832_tp2_sp2|cross/commit/9|0.0495|0.0603|0.0492|0.0870|0.0482|0.6188|0.1097|
|480x832_tp2_sp2|cross/commit/27|0.0494|0.0601|0.0493|0.0867|0.0468|0.6190|0.1074|
|480x832_tp2_sp2|cross/denoise/0|0.0494|0.0599|0.0492|0.0870|0.0471|0.6187|0.1100|
|480x832_tp2_sp2|cross/denoise/9|0.0488|0.0599|0.0495|0.0869|0.0472|0.6187|0.1094|
|480x832_tp2_sp2|cross/denoise/27|0.0493|0.0593|0.0493|0.0869|0.0491|0.6186|0.1087|
|480x832_tp2_sp2|self/commit/0|0.3322|0.1835|0.3576|0.6992|0.2051|4.8254|0.1695|
|480x832_tp2_sp2|self/commit/9|1.2801|0.6898|1.3796|2.7441|0.8088|18.9312|0.6824|
|480x832_tp2_sp2|self/commit/27|1.9207|1.0289|2.0734|4.1031|1.2140|28.3940|1.0151|
|480x832_tp2_sp2|self/denoise/0|0.3334|0.1886|0.3579|0.7013|0.2054|4.8422|0.1750|
|480x832_tp2_sp2|self/denoise/9|1.2806|0.6922|1.3852|2.7511|0.8210|18.9206|0.6898|
|480x832_tp2_sp2|self/denoise/27|1.9247|1.0457|2.0810|4.1043|1.2238|28.3963|1.0237|
|480x832_tp4_sp1|cross/commit/0|0.0514|0.0634|0.0492|0.0867|0.0496|0.6166|0.1111|
|480x832_tp4_sp1|cross/commit/9|0.0463|0.0639|0.0498|0.0868|0.0491|0.6180|0.1104|
|480x832_tp4_sp1|cross/commit/27|0.0514|0.0630|0.0492|0.0868|0.0494|0.6175|0.1113|
|480x832_tp4_sp1|cross/denoise/0|0.0513|0.0619|0.0492|0.0867|0.0490|0.6173|0.1093|
|480x832_tp4_sp1|cross/denoise/9|0.0512|0.0627|0.0493|0.0870|0.0485|0.6166|0.1509|
|480x832_tp4_sp1|cross/denoise/27|0.0516|0.0623|0.0492|0.0867|0.0489|0.6181|0.1093|
|480x832_tp4_sp1|self/commit/0|0.3307|0.1816|0.3569|0.7020|0.2038|4.8002|0.1651|
|480x832_tp4_sp1|self/commit/9|1.2685|0.6716|1.3742|2.7463|0.8023|18.7392|0.6693|
|480x832_tp4_sp1|self/commit/27|1.9001|1.0191|2.0584|4.1153|1.1979|28.1647|1.0006|
|480x832_tp4_sp1|self/denoise/0|0.3316|0.1820|0.3569|0.7018|0.2071|4.8008|0.1737|
|480x832_tp4_sp1|self/denoise/9|1.2705|0.6791|1.3796|2.7516|0.8054|18.7683|0.6761|
|480x832_tp4_sp1|self/denoise/27|1.8932|1.0251|2.0552|4.1194|1.1942|28.0226|1.0042|
|704x1280_tp1_sp1|cross/commit/0|0.3246|0.2654|0.3466|0.6602|0.1954|4.6512|0.1882|
|704x1280_tp1_sp1|cross/commit/9|0.3247|0.2654|0.3468|0.6602|0.1957|4.6520|0.1871|
|704x1280_tp1_sp1|cross/commit/27|0.3258|0.2661|0.3467|0.6610|0.1947|4.6539|0.1875|
|704x1280_tp1_sp1|cross/denoise/0|0.3248|0.2651|0.3467|0.6604|0.1966|4.6513|0.1891|
|704x1280_tp1_sp1|cross/denoise/9|0.3251|0.2648|0.3470|0.6609|0.1945|4.6542|0.1878|
|704x1280_tp1_sp1|cross/denoise/27|0.3248|0.2661|0.3470|0.6606|0.1956|4.6530|0.1885|
|704x1280_tp1_sp1|self/commit/0|6.0150|3.4337|6.5341|12.5485|4.0311|88.2373|3.3911|
|704x1280_tp1_sp1|self/commit/9|24.0640|13.4342|25.7906|50.1805|14.8068|oom|13.4560|
|704x1280_tp1_sp1|self/commit/27|36.1083|20.0373|38.8496|75.2259|22.3491|oom|20.0399|
|704x1280_tp1_sp1|self/denoise/0|6.0974|3.5466|6.5835|12.6717|4.1053|88.2268|3.4814|
|704x1280_tp1_sp1|self/denoise/9|24.1622|13.5304|25.8526|50.3580|14.9247|oom|13.5762|
|704x1280_tp1_sp1|self/denoise/27|36.2916|20.0324|38.7065|75.4718|22.3066|oom|20.2067|
|704x1280_tp1_sp2|cross/commit/0|0.1696|0.1421|0.1809|0.3356|0.1031|2.4256|0.1226|
|704x1280_tp1_sp2|cross/commit/9|0.1707|0.1424|0.1825|0.3364|0.1034|2.4257|0.1273|
|704x1280_tp1_sp2|cross/commit/27|0.1699|0.1418|0.1801|0.3360|0.1031|2.4274|0.1232|
|704x1280_tp1_sp2|cross/denoise/0|0.1696|0.1430|0.1814|0.3365|0.1030|2.4265|0.1261|
|704x1280_tp1_sp2|cross/denoise/9|0.1696|0.1419|0.1808|0.3363|0.1030|2.4271|0.1264|
|704x1280_tp1_sp2|cross/denoise/27|0.1699|0.1425|0.1811|0.3364|0.1033|2.4261|0.1253|
|704x1280_tp1_sp2|self/commit/0|3.0658|1.7203|3.3167|6.3527|2.0383|44.9385|1.6880|
|704x1280_tp1_sp2|self/commit/9|12.2549|6.8533|13.2156|25.2085|7.5811|187.5624|6.8607|
|704x1280_tp1_sp2|self/commit/27|18.4225|10.2620|19.9064|37.8296|11.3507|281.0248|10.2877|
|704x1280_tp1_sp2|self/denoise/0|3.0994|1.7655|3.3569|6.4097|2.0885|44.9332|1.7614|
|704x1280_tp1_sp2|self/denoise/9|12.3170|6.8844|13.2669|25.2830|7.6094|187.7495|6.9286|
|704x1280_tp1_sp2|self/denoise/27|18.4596|10.2998|19.9718|37.8863|11.3978|280.8996|10.3469|
|704x1280_tp1_sp4|cross/commit/0|0.0962|0.0812|0.1029|0.1793|0.0596|1.2663|0.1259|
|704x1280_tp1_sp4|cross/commit/9|0.0961|0.0814|0.1028|0.1791|0.0596|1.2643|0.1266|
|704x1280_tp1_sp4|cross/commit/27|0.0973|0.0809|0.1029|0.1791|0.0597|1.2651|0.1248|
|704x1280_tp1_sp4|cross/denoise/0|0.0959|0.0803|0.1030|0.1792|0.0597|1.2654|0.1255|
|704x1280_tp1_sp4|cross/denoise/9|0.0971|0.0809|0.1029|0.1792|0.0602|1.2670|0.1263|
|704x1280_tp1_sp4|cross/denoise/27|0.0973|0.0808|0.1029|0.1795|0.0600|1.2669|0.1282|
|704x1280_tp1_sp4|self/commit/0|1.6080|0.8818|1.7358|3.3039|1.0398|23.3521|0.8616|
|704x1280_tp1_sp4|self/commit/9|6.3741|3.5161|6.9239|13.2254|3.8682|97.4513|3.4944|
|704x1280_tp1_sp4|self/commit/27|9.6274|5.2461|10.3836|19.8428|5.8081|146.3963|5.2219|
|704x1280_tp1_sp4|self/denoise/0|1.6290|0.9015|1.7560|3.3056|1.0537|23.3936|0.8846|
|704x1280_tp1_sp4|self/denoise/9|6.4125|3.5336|6.9458|13.2361|3.8888|97.4931|3.5063|
|704x1280_tp1_sp4|self/denoise/27|9.6641|5.2757|10.4008|19.8536|5.8391|146.4556|5.2401|
|704x1280_tp2_sp1|cross/commit/0|0.1693|0.1425|0.1807|0.3360|0.1029|2.4241|0.1222|
|704x1280_tp2_sp1|cross/commit/9|0.1689|0.1416|0.1807|0.3365|0.1023|2.4240|0.1328|
|704x1280_tp2_sp1|cross/commit/27|0.1691|0.1422|0.1821|0.3387|0.1025|2.4251|0.1244|
|704x1280_tp2_sp1|cross/denoise/0|0.1691|0.1420|0.1808|0.3360|0.1023|2.4246|0.1231|
|704x1280_tp2_sp1|cross/denoise/9|0.1690|0.1427|0.1807|0.3362|0.1021|2.4246|0.1238|
|704x1280_tp2_sp1|cross/denoise/27|0.1691|0.1425|0.1809|0.3366|0.1028|2.4243|0.1225|
|704x1280_tp2_sp1|self/commit/0|3.0029|1.7114|3.2718|6.3710|1.9977|44.9897|1.6797|
|704x1280_tp2_sp1|self/commit/9|12.1157|6.7450|13.0450|24.9920|7.4169|186.3893|6.7404|
|704x1280_tp2_sp1|self/commit/27|18.1780|10.1017|19.5554|37.8624|11.1017|279.7005|10.0892|
|704x1280_tp2_sp1|self/denoise/0|3.0552|1.7263|3.3071|6.3689|2.0448|45.0016|1.7069|
|704x1280_tp2_sp1|self/denoise/9|12.1563|6.7615|13.0861|25.0362|7.4504|186.4961|6.7833|
|704x1280_tp2_sp1|self/denoise/27|18.2043|10.1238|19.6090|37.5706|11.1277|279.7636|10.1260|
|704x1280_tp2_sp2|cross/commit/0|0.0974|0.0816|0.1038|0.1795|0.0600|1.2690|0.1329|
|704x1280_tp2_sp2|cross/commit/9|0.0972|0.0823|0.1035|0.1796|0.0601|1.2699|0.1208|
|704x1280_tp2_sp2|cross/commit/27|0.0974|0.0812|0.1034|0.1797|0.0607|1.2706|0.1214|
|704x1280_tp2_sp2|cross/denoise/0|0.0972|0.0807|0.1033|0.1796|0.0600|1.2703|0.1227|
|704x1280_tp2_sp2|cross/denoise/9|0.0971|0.0812|0.1031|0.1796|0.0601|1.2728|0.1262|
|704x1280_tp2_sp2|cross/denoise/27|0.0975|0.0814|0.1036|0.1798|0.0599|1.2692|0.1223|
|704x1280_tp2_sp2|self/commit/0|1.6325|0.8971|1.7428|3.3378|1.0528|23.5017|0.8731|
|704x1280_tp2_sp2|self/commit/9|6.4719|3.5630|7.0021|13.2380|3.9482|98.4058|3.5638|
|704x1280_tp2_sp2|self/commit/27|9.7437|5.3452|10.4907|19.8512|5.9283|147.7506|5.3355|
|704x1280_tp2_sp2|self/denoise/0|1.6512|0.9164|1.7573|3.3401|1.0693|23.5813|0.8995|
|704x1280_tp2_sp2|self/denoise/9|6.5002|3.5929|7.0194|13.2475|3.9660|98.3540|3.5810|
|704x1280_tp2_sp2|self/denoise/27|9.8104|5.3796|10.5177|19.8653|5.9443|147.7387|5.3451|
|704x1280_tp4_sp1|cross/commit/0|0.0973|0.0805|0.1028|0.1791|0.0598|1.2653|0.1247|
|704x1280_tp4_sp1|cross/commit/9|0.0968|0.0813|0.1028|0.1792|0.0596|1.2648|0.1257|
|704x1280_tp4_sp1|cross/commit/27|0.0960|0.0807|0.1026|0.1790|0.0598|1.2648|0.1248|
|704x1280_tp4_sp1|cross/denoise/0|0.0964|0.0809|0.1029|0.1790|0.0596|1.2650|0.1244|
|704x1280_tp4_sp1|cross/denoise/9|0.0962|0.0806|0.1028|0.1793|0.0595|1.2654|0.1244|
|704x1280_tp4_sp1|cross/denoise/27|0.0960|0.0806|0.1029|0.1791|0.0594|1.2658|0.1267|
|704x1280_tp4_sp1|self/commit/0|1.5977|0.8807|1.7306|3.3322|1.0424|23.3082|0.8584|
|704x1280_tp4_sp1|self/commit/9|6.4078|3.5067|6.9284|13.2265|3.8713|97.4816|3.4978|
|704x1280_tp4_sp1|self/commit/27|9.6880|5.2711|10.3760|19.8481|5.8084|146.2959|5.2081|
|704x1280_tp4_sp1|self/denoise/0|1.6306|0.8934|1.7469|3.3057|1.0619|23.3940|0.8837|
|704x1280_tp4_sp1|self/denoise/9|6.4128|3.5200|6.9477|13.2383|3.8842|97.4906|3.5117|
|704x1280_tp4_sp1|self/denoise/27|9.6495|5.2852|10.4068|19.8566|5.8207|146.3604|5.2382|

### 已审核点的数值误差

下表为各后端在已审核输入中的最大观测值，各列最大值可能来自不同输入。抽样FP64参考只覆盖至多32个query位置、4个head及全部有效K/V，不能视为完整输出误差上界。完整原始输出另作对照；预设relative L2门槛为0.01，allclose使用atol=rtol=0.03。

|后端|审核点数|对保存输出最大relative L2|对FP64抽样最大relative L2|对FP64抽样最大绝对误差|
|---|---:|---:|---:|---:|
|native_dense|144|0.000921|0.002221|0.016575|
|sdpa_cudnn|144|0.001503|0.002231|0.015672|
|sdpa_efficient|144|0.001506|0.002231|0.015672|
|sdpa_flash|144|0.001011|0.002220|0.016575|
|sdpa_math|140|0.002165|0.001750|0.015578|
|vllm_fa2|144|0.001010|0.002220|0.016575|
|vllm_fa3|144|0.000326|0.002220|0.016575|

[下载逐点误差、输入标识及审核路径 CSV](/attention_errors.csv)。这些是孤立Attention结果，不证明整模型数值等价。


### 已核验配置的可下载图与数据

![480x832_tp1_sp1 Attention 实测散点图](/attention_480x832_tp1_sp1.png)

[下载 480x832_tp1_sp1 矢量图 SVG](/attention_480x832_tp1_sp1.svg)

![480x832_tp1_sp1 同输入耗时比值热力图](/speedup_480x832_tp1_sp1.png)

![480x832_tp1_sp2 Attention 实测散点图](/attention_480x832_tp1_sp2.png)

[下载 480x832_tp1_sp2 矢量图 SVG](/attention_480x832_tp1_sp2.svg)

![480x832_tp1_sp2 同输入耗时比值热力图](/speedup_480x832_tp1_sp2.png)

![480x832_tp1_sp4 Attention 实测散点图](/attention_480x832_tp1_sp4.png)

[下载 480x832_tp1_sp4 矢量图 SVG](/attention_480x832_tp1_sp4.svg)

![480x832_tp1_sp4 同输入耗时比值热力图](/speedup_480x832_tp1_sp4.png)

![480x832_tp2_sp1 Attention 实测散点图](/attention_480x832_tp2_sp1.png)

[下载 480x832_tp2_sp1 矢量图 SVG](/attention_480x832_tp2_sp1.svg)

![480x832_tp2_sp1 同输入耗时比值热力图](/speedup_480x832_tp2_sp1.png)

![480x832_tp2_sp2 Attention 实测散点图](/attention_480x832_tp2_sp2.png)

[下载 480x832_tp2_sp2 矢量图 SVG](/attention_480x832_tp2_sp2.svg)

![480x832_tp2_sp2 同输入耗时比值热力图](/speedup_480x832_tp2_sp2.png)

![480x832_tp4_sp1 Attention 实测散点图](/attention_480x832_tp4_sp1.png)

[下载 480x832_tp4_sp1 矢量图 SVG](/attention_480x832_tp4_sp1.svg)

![480x832_tp4_sp1 同输入耗时比值热力图](/speedup_480x832_tp4_sp1.png)

![704x1280_tp1_sp1 Attention 实测散点图](/attention_704x1280_tp1_sp1.png)

[下载 704x1280_tp1_sp1 矢量图 SVG](/attention_704x1280_tp1_sp1.svg)

![704x1280_tp1_sp1 同输入耗时比值热力图](/speedup_704x1280_tp1_sp1.png)

![704x1280_tp1_sp2 Attention 实测散点图](/attention_704x1280_tp1_sp2.png)

[下载 704x1280_tp1_sp2 矢量图 SVG](/attention_704x1280_tp1_sp2.svg)

![704x1280_tp1_sp2 同输入耗时比值热力图](/speedup_704x1280_tp1_sp2.png)

![704x1280_tp1_sp4 Attention 实测散点图](/attention_704x1280_tp1_sp4.png)

[下载 704x1280_tp1_sp4 矢量图 SVG](/attention_704x1280_tp1_sp4.svg)

![704x1280_tp1_sp4 同输入耗时比值热力图](/speedup_704x1280_tp1_sp4.png)

![704x1280_tp2_sp1 Attention 实测散点图](/attention_704x1280_tp2_sp1.png)

[下载 704x1280_tp2_sp1 矢量图 SVG](/attention_704x1280_tp2_sp1.svg)

![704x1280_tp2_sp1 同输入耗时比值热力图](/speedup_704x1280_tp2_sp1.png)

![704x1280_tp2_sp2 Attention 实测散点图](/attention_704x1280_tp2_sp2.png)

[下载 704x1280_tp2_sp2 矢量图 SVG](/attention_704x1280_tp2_sp2.svg)

![704x1280_tp2_sp2 同输入耗时比值热力图](/speedup_704x1280_tp2_sp2.png)

![704x1280_tp4_sp1 Attention 实测散点图](/attention_704x1280_tp4_sp1.png)

[下载 704x1280_tp4_sp1 矢量图 SVG](/attention_704x1280_tp4_sp1.svg)

![704x1280_tp4_sp1 同输入耗时比值热力图](/speedup_704x1280_tp4_sp1.png)

![mfu_480x832_tp1_sp1 Attention 实测散点图](/attention_mfu_480x832_tp1_sp1.png)

[下载 mfu_480x832_tp1_sp1 矢量图 SVG](/attention_mfu_480x832_tp1_sp1.svg)

![mfu_480x832_tp1_sp2 Attention 实测散点图](/attention_mfu_480x832_tp1_sp2.png)

[下载 mfu_480x832_tp1_sp2 矢量图 SVG](/attention_mfu_480x832_tp1_sp2.svg)

![mfu_480x832_tp1_sp4 Attention 实测散点图](/attention_mfu_480x832_tp1_sp4.png)

[下载 mfu_480x832_tp1_sp4 矢量图 SVG](/attention_mfu_480x832_tp1_sp4.svg)

![mfu_480x832_tp2_sp1 Attention 实测散点图](/attention_mfu_480x832_tp2_sp1.png)

[下载 mfu_480x832_tp2_sp1 矢量图 SVG](/attention_mfu_480x832_tp2_sp1.svg)

![mfu_480x832_tp2_sp2 Attention 实测散点图](/attention_mfu_480x832_tp2_sp2.png)

[下载 mfu_480x832_tp2_sp2 矢量图 SVG](/attention_mfu_480x832_tp2_sp2.svg)

![mfu_480x832_tp4_sp1 Attention 实测散点图](/attention_mfu_480x832_tp4_sp1.png)

[下载 mfu_480x832_tp4_sp1 矢量图 SVG](/attention_mfu_480x832_tp4_sp1.svg)

![mfu_704x1280_tp1_sp1 Attention 实测散点图](/attention_mfu_704x1280_tp1_sp1.png)

[下载 mfu_704x1280_tp1_sp1 矢量图 SVG](/attention_mfu_704x1280_tp1_sp1.svg)

![mfu_704x1280_tp1_sp2 Attention 实测散点图](/attention_mfu_704x1280_tp1_sp2.png)

[下载 mfu_704x1280_tp1_sp2 矢量图 SVG](/attention_mfu_704x1280_tp1_sp2.svg)

![mfu_704x1280_tp1_sp4 Attention 实测散点图](/attention_mfu_704x1280_tp1_sp4.png)

[下载 mfu_704x1280_tp1_sp4 矢量图 SVG](/attention_mfu_704x1280_tp1_sp4.svg)

![mfu_704x1280_tp2_sp1 Attention 实测散点图](/attention_mfu_704x1280_tp2_sp1.png)

[下载 mfu_704x1280_tp2_sp1 矢量图 SVG](/attention_mfu_704x1280_tp2_sp1.svg)

![mfu_704x1280_tp2_sp2 Attention 实测散点图](/attention_mfu_704x1280_tp2_sp2.png)

[下载 mfu_704x1280_tp2_sp2 矢量图 SVG](/attention_mfu_704x1280_tp2_sp2.svg)

![mfu_704x1280_tp4_sp1 Attention 实测散点图](/attention_mfu_704x1280_tp4_sp1.png)

[下载 mfu_704x1280_tp4_sp1 矢量图 SVG](/attention_mfu_704x1280_tp4_sp1.svg)

![mfu_overview Attention 实测散点图](/attention_mfu_overview.png)

[下载 mfu_overview 矢量图 SVG](/attention_mfu_overview.svg)

![winners Attention 实测散点图](/attention_winners.png)

[下载 winners 矢量图 SVG](/attention_winners.svg)

图中点是三组均值的中位数，误差棒为三组的最小/最大值，不是置信区间。横坐标为离散输入组；未连接或插值缺失点。self与cross的后端排序不同，不能相互外推。

[下载全部逐组图表数据 CSV](/attention_graph_points.csv)

热力图仅计算同一保存输入的耗时比值：self参考FA3 varlen，cross参考native dense；大于1表示本次孤立回放中比参考耗时更短。缺少参考的格子保持missing。比值使用三组均值的中位数，不是统计显著性或端到端加速证明。颜色范围截断时，格内数值仍为完整实测比值。

[下载比值及参考后端 JSON](/attention_speedup_points.json)

归档索引扫描时间：2026-09-10T08:50:17.874989+00:00；记录37069个文件、约407.04GB。这是扫描时点快照，活动文件仍需在任务结束后重新封存。索引位于analysis/raw_data_index.json。

### FA3接口差异的直接证据

已逐点核对144组相同输入的native dense与FA3 varlen实际dispatch。以下为已审核输入的观察范围，不是置信区间。

|角色|输入对数|两者均为SM90 FA3|主kernel名称不同|native/varlen耗时比范围|
|---|---:|---:|---:|---:|
|self|72|72|72|0.909–1.010|
|cross|72|72|72|0.705–2.407|

144/144组varlen接口观察到prepare_varlen准备kernel。不能把耗时差异全归因于准备开销，也不能将native dense更慢误判为FA2 fallback。具体调度与资源利用原因仍需硬件计数器和单变量实验。

完整kernel名称、输入路径、计时与审核链接见analysis/fa3_interface_dispatch_comparison.json；复现脚本为scripts/compare_fa3_interfaces.py。


现有FA2/FA3数据使用varlen接口；原模型cross默认路径使用普通dense接口，两者开销可能不同。后续完整矩阵已包含native dense对照，详见逐点接口审核；整模型cross候选未通过数值验收，不能认领替换收益。

## 硬件指标：首次真实FA3输入试采

输入来自480×832 TP1/SP4的self-attention，commit、起始帧27；Q=[4680,10,128]，K存储=[29640,10,128]，有效K长度28080。root在普通用户gpu run分配的一张卡内采集。总任务约20.7秒，两个kernel各11轮replay。

主Attention kernel的NCU时间923.072µs，SM throughput为工具定义峰值的79.48%，DRAM throughput 4.44%，L2 throughput 23.48%，achieved occupancy 18.75%；每线程168寄存器，每block分配约200.704KB shared memory。
这些值属于一次隔离replay及NCU所定义的指标，不是模型MFU。DRAM吞吐低、SM吞吐较高提供优先检查计算/执行资源的线索；尚无反证实验，不能据此确认唯一瓶颈。低occupancy本身也不证明需要提高occupancy。无profiler的回放计时与该NCU时间分别报告。

原始证据：raw/ncu_attention_fa3_pilot_v1/attention.ncu-rep、metrics.csv、selected_metrics.json；原始版本/命令/输入哈希见manifest.json。

## 未完成项目及影响

新增限制：后续FA2 NCU试采被调度守护程序终止；普通用户分配下的root进程被判为foreign，--user root标签也未解决。早先已保存的FA3计数器仍是有效实测，但不能推断后续root任务可正常长时间执行。当前暂停此类重试，普通用户调度要求保持不变；需要可兼容该要求的root采集权限方案才能继续。未改用root调度，未修改守护或驱动。

|项目|状态及影响|
|---|---|
|两种分辨率、6种TP/SP配置|纯DiT输入矩阵12/12已覆盖；1008项回放中1004项有效、4项SDPA math OOM。高分辨率TP1/SP1完整VAE流水线采集失败，不能将纯DiT覆盖等同完整流水线成功|
|TRTLLM_ATTN兼容性候选|当前安装路径源码只接受SM100/103/107，实测设备SM90；静态判定不兼容，未执行GPU试跑、无耗时数值。另在采集启动日志观察到FlashInfer 0.6.16.post3与flashinfer-cubin 0.6.13不匹配的可选导入警告；这不证明实际模型选择的后端发生fallback。证据：analysis/trtllm_compatibility.json、analysis/flashinfer_import_warning.json；不包含在7接口回放分母中|
|QKV all-to-all前后及paged/gather完整关联|现有代表张量不等于完整前后链路；关联审计未完成|
|H0参考优化适用性|已保存H1–H5本地源码映射；编译后融合与完整transport关联仍缺失，所列外部参考状态已追加核验，完整运行适用性仍未通过|
|H1–H5优化消融|未完成；当前uint8输出已存在，不能重复认领收益；H3专用VAE不能直接替换FP32 Wan VAE|
|全部层与shape组映射|未完成；首层代表输入不等于完整模型shape覆盖|
|矩阵乘、VAE卷积、数据转换计数器|66条真实输入已保存，18项rank0回放通过逐位检查；硬件计数器仍未完成，不能宣称全流水线瓶颈归因完成|
|硬件峰值校准、MFU、DRAM/L2 Roofline|未完成；不能把SM throughput百分比当MFU|
|优选后端整模型A/B|已执行cuDNN cross候选，包括eager控制；40份候选输出均未通过预设数值门槛，尚无通过验收的整模型优化|
|最终独立重复与尾延迟|两种分辨率各三次、每次1000稳态间隔均完成并通过计时复核；未设跨启动稳定性通过阈值，不能保证未来尾延迟，输出正确性门槛仍未通过|
|编码、网络、背压和取消|未测；CPU-ready FPS不代表服务SLA|

## 参考优化源码核验

analysis/optimization_transfer_matrix.json记录已检查文件的SHA256、符号、必要条件、已启用状态与拒绝原因。当前H0仍为部分归因：LingBot已融合QKV投影，并在Ulysses>1时执行5D合并QKV交换；实际metadata确认前后shape，不能重复认领合并collective收益。当前路径含layout materialization并调用dist.all_to_all_single，不能称native SymmMem已启用。TP RMSNorm保留全局平方和归约，融合必须保持语义；编译后具体融合仍待关联。H3 VAE与Wan接口不同，GPU uint8转换已存在，Case D顺序decode。未把源码存在或参考提案视作性能收益。

## 原始数据与复现

持久目录：/data/zhoutaichang/feature/Lingbot_world_realtime/runs/analysis_report_20260908。新增实验使用独立目录，不覆盖旧原始数据。raw/evening_matrix_v1.json保存顺序任务命令和结果；每次capture的manifest保存冻结脚本哈希、输入配置、时间和分配GPU。每次replay保存results.json、timings.jsonl及每个后端实际dispatch trace。analysis下audit文件链接输入校验和及计时验证。完整归档校验清单仍需补齐。
脚本位于scripts；必须在gpu run分配范围内运行GPU任务，root仅用于所需的NCU采集。不使用gpu reserve，不修改驱动或其他用户进程。

## 逐kernel待归因清单

v2完整时间线的1091749次记录调用已按完整kernel名称汇总；稳态有102个名称分组。analysis/bottlenecks.json逐组列出尚缺的实际shape/阶段、匹配硬件计数器及反证实验，当前均为unresolved。名称分组不等于shape分组；时钟与末块存在性通过不等于逐层调用完整，也不意味着瓶颈已确定。旧记录的限制独立保留。

[下载调用覆盖清单](/kernel-coverage.json) · [下载逐组待归因清单](/kernel-bottlenecks.json)。生成脚本：scripts/build_kernel_coverage.py。
## 最新补采：同名kernel的启动参数差异

稳态窗口内的176,984次调用按物理GPU、完整kernel名称、grid/block、寄存器和共享内存分成970组。102个kernel名称中，36个名称对应多种启动参数组合（该比较排除GPU编号）。调用数与累计GPU时间均与原有稳态排名复核一致。

这说明仅按名称合并会掩盖启动配置差异；后续NCU采样需保留对应参数及真实输入。启动参数仍不等于张量shape，也不提供实际occupancy、带宽或warp stall。跨GPU累计时间不等于墙钟时间。

[下载按物理GPU划分的完整启动参数CSV](/worker-launch-resources-steady.csv)。证据：analysis/worker_launch_resources_summary.json；复算脚本：scripts/summarize_worker_launch_resources.py。


## 交付源文件快照与归档核对

源文件快照：raw/delivery_source_snapshot_20260910T002011Z，包含116个分析/发布脚本、报告与索引文件，各文件复制前后SHA256一致。原始数据仍保留在各独立运行目录。

2026-09-10T00:20:42.953404+00:00核对原始索引中的30,217个文件（约326.42GB），文件存在性、大小和mtime异常共0项。该检查使用增量哈希索引，未重新读取全部326GB内容计算哈希，不等于最终不可变封存。

历史未保存的完整输出仍为缺失。索引只能证明列入的文件，不能证明所有计划数据均已采集。审核记录：analysis/delivery_archive_audit_20260910T0020Z.json；源快照内manifest.json保存逐文件来源与校验和。


## 数值差异从何时出现

三组未通过的比较（reduce-overhead重复、default重复、eager cuDNN候选）在两轮的首个chunk均已未通过完整数值门槛。误差不是只在长缓存或末块出现。末块latent相对L2较首块更大，但中间存在回落，不能称单调累积，也不能仅凭此曲线证明缓存、编译器或某个kernel是唯一原因。eager重复的全部已保存输出逐位一致。

![逐chunk输出误差](/execution_error_progression.png)

|比较（epoch 1）|首块latent相对L2|末块latent相对L2|通过输出数（latent+像素）|
|---|---:|---:|---:|
|eager repeat|0.000%|0.000%|20/20|
|compiled reduce-overhead repeat|1.980%|24.126%|0/20|
|compiled default repeat|1.286%|21.170%|0/20|
|eager cuDNN candidate vs eager reference|3.445%|26.139%|0/20|

图中的虚线只表示所绘指标的门槛分量，不代表完整验收：像素RMSE低于0.01仍可能因max_abs超过0.03失败；latent还要求allclose。像素按[-1,1]归一化，右侧纵轴为RMSE乘100，不是感知质量损失百分比。每条曲线只有一对启动、每轮十个chunk，没有总体置信区间。

[下载逐点CSV](/execution_error_progression.csv) · [下载矢量SVG](/execution_error_progression.svg)。复算脚本：scripts/plot_execution_error_progression.py；数据及来源SHA256：analysis/execution_error_progression.json。


## 按交付间隔检查四卡活动

最新worker写出补采的四个稳态交付间隔逐段计算区间并集与交集，再按物理GPU核对。非NCCL、NCCL及两者重叠的分段总量均与已有整段统计一致。下表为四张卡之间的实测最小–最大范围，不是置信区间。

|交付chunk|窗口ms|非NCCL并集ms（四卡范围）|NCCL并集ms（四卡范围）|两者重叠ms（四卡范围）|最长无kernel记录区间ms（四卡范围）|
|---|---:|---:|---:|---:|---:|
|5|847.494|672.786–681.350|72.699–88.814|15.424–21.212|40.557–40.806|
|6|853.145|670.476–677.783|74.130–84.899|16.278–21.249|48.950–49.153|
|7|853.728|669.866–679.622|68.339–80.790|15.369–20.147|47.190–47.676|
|8|856.184|670.028–682.199|62.485–92.971|14.744–21.930|49.731–50.412|

非NCCL并集约670–682ms，但不能称全部为计算；其中也有重排、转换等kernel。NCCL活动存在卡间差异，尚不能证明哪张卡是请求关键路径上的慢卡。活动重叠不证明硬件并发执行或收益。

最长无kernel记录区间约41–50ms，仅由kernel表求补集，未扣除memcpy、memset或CPU活动，不能当作GPU空闲、CPU等待或可直接消除的开销。当前仅四个profiler下间隔，不用于尾延迟验收。

[下载16行逐卡逐chunk CSV](/worker_chunk_activity.csv)。完整数据：analysis/worker_chunk_activity.json；复算脚本：scripts/analyze_worker_chunk_activity.py。


## 最长无kernel记录区间：拷贝与主机交接核对

在16个逐卡逐chunk最长区间中，同一物理GPU的memcpy/memset仅覆盖约0.0007–0.0044ms，不能解释约41–50ms的主体。进程内CUDA runtime调用也已保留，但它们只是上下文关联；多GPU进程中的API不能自动归属于这张卡，累计线程耗时不能当作墙钟等待。

|chunk|前块CPU-ready→结果收到ms|结果收到→下一块DiT调用ms|
|---|---:|---:|
|5|16.286|2.556|
|6|22.506|2.621|
|7|22.337|2.542|
|8|24.322|3.981|

这些最长区间位于前块CPU-ready之后、下一块开始交接附近。冻结worker源码在记录ready后执行finite检查、按条件保留/保存输出、计算像素SHA256，再发送结果；主进程收到结果后记录日志并提交下一块事件。这些步骤可能贡献块间间隔，但没有逐操作计时，不能将16–24ms全部归因于哈希、IPC或调度。

这也限定了CPU-ready交付口径：逐块ready间隔会包含前一块ready之后的基准记录与交接工作；它不是只算GPU核心执行的时间。当前结果来自profiler补采，不能将这段间隔直接从未插桩FPS中扣除或认领服务优化收益。要量化贡献，需在相同输出协议下单独计时或做受控对照。

[下载区间边界、拷贝事件及主机API上下文JSON](/worker_longest_gap_context.json)。原始SQLite事件rowid、物理映射、冻结源码SHA256与主机时钟标记均保存在该文件。复算脚本：scripts/analyze_worker_longest_gaps.py。


## 无profiler扩样：交接时间与最慢样本

复算六次正常计时运行的6,000个稳态间隔，每次250个session、每个session四个间隔。这里的交接定义为“前块CPU-ready到下一块DiT调用”，涵盖校验、结果传递、日志和调度等路径；不是某个单独函数耗时。每个样本均验证：交接时间 + 下一块DiT调用至当前CPU-ready = 当前交付间隔。

|运行|交接均值ms|P50 ms|P99 ms|最大ms|占平均交付间隔|
|---|---:|---:|---:|---:|---:|
|tail_original_480x832_r0_v1|16.171|15.676|22.198|212.028|1.96%|
|tail_original_480x832_r1_v1|17.349|15.863|29.812|97.176|2.08%|
|tail_original_480x832_r2_v1|16.042|15.584|26.293|58.167|1.93%|
|tail_original_704x1280_r0_v1|32.167|30.653|82.375|207.263|1.31%|
|tail_original_704x1280_r1_v1|32.086|30.681|60.496|194.406|1.30%|
|tail_original_704x1280_r2_v1|32.468|30.765|46.808|51.999|1.32%|

|运行的最慢交付样本|epoch / chunk|完整间隔ms|其中交接ms|
|---|---|---:|---:|
|tail_original_480x832_r0_v1|129 / 7|1028.816|212.028|
|tail_original_480x832_r1_v1|32 / 8|1010.481|17.592|
|tail_original_480x832_r2_v1|162 / 6|881.423|17.885|
|tail_original_704x1280_r0_v1|217 / 6|2899.883|119.889|
|tail_original_704x1280_r1_v1|11 / 7|2980.876|111.208|
|tail_original_704x1280_r2_v1|163 / 5|2711.041|30.333|

480×832 r0的最慢交付样本含约212ms交接，而r1、r2最慢样本的交接仅约18ms。不同长尾样本落在不同阶段，不能统一归因于GPU计算、IPC或哈希。以上是记录边界的分解，不是证明删除交接会得到等额收益；没有做相应受控优化。

P99采用排序样本线性插值，仅为各运行描述统计，session内样本相关，不当作1,000次独立启动。分量均值可相加，分量P99不可相加。无profiler仍保留正常基准记录和输出协议，不能称为无任何测量开销。

[下载6,000行逐次交接与交付时间CSV](/clean_handoff_intervals.csv) · [下载分组统计及最慢样本JSON](/clean_handoff_summary.json)。复算脚本：scripts/analyze_clean_handoff.py；来源summary的SHA256保存在统计文件中。


## 算子回放与模型调用：启动参数匹配的歧义

18个已审核rank0算子回放包含35种完整名称/启动参数组合；其中10种在最新worker补采的稳态模型记录中匹配，3种匹配组合对应多个已保存输入。这里分母是启动参数组合，不是算子输入数或模型shape数。

匹配字段包括完整kernel名称、grid、block、寄存器和静态加动态共享内存字节数。3组歧义分别涉及GEMM、另一GEMM及VAE卷积；输入身份不同不一定意味着数学shape不同。一个GEMM组合同时对应输出[4680,5120]和[1,4680,5120]，单例维度也可能对应相同展平矩阵，不能借此声称两种计算量不同。

即使只有一个已保存候选，未捕获输入仍可能使用同一启动组合，故全部对应关系继续标为unresolved。未匹配也不能证明模型没有该算子，可能涉及编译、算法选择或窗口差异。这一步缩小后续采样范围，但不完成K1逐层shape归因，也不能将孤立回放计数器直接赋给模型全部同名调用。

[下载候选输入、启动参数与模型事件ID](/operator_signature_ambiguity.json)。复算脚本：scripts/audit_operator_signature_ambiguity.py。


## 首版之后补充：H3外部参考状态

核验时间：2026-09-10T01:42:08.017484+00:00（UTC）。此补充晚于08:40固定报告包，在线报告追加更新，不修改原包。

|参考|外部状态|对LingBot分析的含义|
|---|---|---|
|[Fast Ulysses #6340](https://github.com/vllm-project/vllm-omni/pull/6340)|2026-08-30已合并，提交8be601ed|可进行源码适用性审查；不能因合并就声称当前LingBot启用SymmMem或取得收益。|
|[H3 VAE算子 #6607](https://github.com/vllm-project/vllm-omni/pull/6607)|2026-08-28已合并，提交04b97a59|H3模型特定算子；不能直接替换FP32 Wan VAE并宣称数值等价。|
|[分块VAE→传输→MP4重叠 RFC #6872](https://github.com/vllm-project/vllm-omni/issues/6872)|核验时仍open|开放提案，不作为LingBot已实现功能或实测加速证据。|

只读本地git核对确认，前两个合并提交均是原始b9cdea7b与rebase c59234d1的祖先；这不排除后续修改或回退，也不能证明相关路径被当前配置执行。未切换checkout、未拉取或修改模型源码。

[H3官方文章](https://vllm.ai/blog/2026-09-01-minimax-h3-production-serving)的第6.3节仍将原始benchmark bundle标为待发布。外部文章数字继续仅作背景，不用来补本机缺失计数器、拟合MFU或承诺端到端收益。

[下载外部响应元数据及SHA256](/h3_external_reference_review.json) · [下载本地提交祖先核对](/h3_reference_local_ancestry.json)。原始网页/API响应保存在raw/h3_external_review_20260910T014207Z。


## 首版之后补充：计数器权限与真实算子补采

2026-09-10 10:00之后，普通用户gpu run分配GPU，SSH root启动后切回实际UID2005，仅本次进程保留CAP_SYS_ADMIN。合成验证采到三个有效计数器、正确性通过、30秒观察完成且无该PID违规记录。无需修改驱动或调度守护；旧root UID执行路线仍停用。此新结果更新前文“等待兼容权限方案”的状态。

三项真实输入试采均成功，原算子与保存输出逐值一致；GEMM、VAE、转换分别耗时约14.15、16.92、11.48秒（包含启动与NCU采集，不是kernel耗时）。后续15个已审核rank0输入也已完成单卡串行补采，汇总见下文。

| 阶段 / kernel | NCU耗时 µs | SM吞吐 % | DRAM吞吐 % | L2吞吐 % | 活跃warp占用 % |
|---|---:|---:|---:|---:|---:|
| dit / nvjet_sm90_tst_256x144_64x4_2x1_v_bz_coopA_bias_TNT | 292.448 | 94.028 | 25.372 | 52.628 | 14.650 |
| vae / void cudnn::nchwToNhwcKernel | 70.336 | 37.540 | 63.710 | 80.816 | 93.308 |
| vae / void cudnn::nchwToNhwcKernel | 3.520 | 12.390 | 5.954 | 15.415 | 24.417 |
| vae / sm90_xmma_fprop_implicit_gemm_f32f32_tf32f32_f32_nhwckrsc_nhwc_tile… | 345.600 | 29.686 | 8.985 | 60.273 | 17.755 |
| vae / void cudnn::nhwcToNchwKernel | 19.904 | 41.754 | 58.781 | 81.369 | 89.588 |
| vae / void at::elementwise_kernel | 27.264 | 59.159 | 39.057 | 60.717 | 79.247 |
| output / void unrolled_elementwise_kernel | 39.584 | 55.888 | 36.344 | 44.802 | 89.899 |

解释：这次GEMM的SM吞吐约94%，不能仅凭14.65%的warp占用判断其低效。VAE卷积本体SM约29.69%、DRAM约8.99%、L2约60.27%，尚不足以确定唯一瓶颈；其输入/输出布局转换L2吞吐约81%，适合作为后续布局对照实验的调查对象。卷积实际kernel名称含TF32，不能因输入FP32而称严格IEEE FP32计算。

边界：表中百分比为NCU相应指标的峰值归一化结果，不是MFU。每个kernel经过11次采集回放；各算子分配到的物理GPU不同，缓存与时钟行为也可能受采集影响。这些是独立算子硬件诊断，不能直接换算整模型加速。完整kernel名称、原生单位、输入/报告SHA256见审核JSON。

[权限验证审核](/capability-counter-probe-audit.json) · [三项真实算子审核与指标](/capability-operator-ncu-audit.json)

### 18个代表输入扩采完成

8个DiT线性算子、8个VAE卷积、2个输出转换全部通过原始输出逐值一致校验，共45次kernel采集（8+35+2）。这些是18个保存输入的独立回放，不是整模型45类kernel的完整覆盖。原生NCU报告、CSV、启动命令、输入SHA256、输出比较、进程身份和对应PID的调度违规查询均已保存。

单项试采总时长为约11.48–21.53秒，含SSH启动、加载、校验及NCU回放；这次数据不支持直接估算全模型NCU耗时。审核文件保留各指标原生单位；例如NCU可能将不同kernel耗时自动显示为ms或µs，跨行计算必须先统一单位。

[18项完整计数器审核与指标](/capability-operator-ncu-full-audit.json)

### 同卡FA2 / FA3硬件对照

输入为480×832 TP1/SP4、frame27 commit self-attention：Q长4680、10头、head_dim128。K/V分配长度29640，但cu_seqlens记录的有效长度为28080。FP64参考严格使用有效K/V，采样32个Q位置、4个头。初次审核误把填充区纳入计算，失败记录与错误脚本保留于raw/capability_attention_pair_oracle_padding_error_v1.*；修正参考后沿用原有0.03/0.03 allclose和1%相对L2门槛，未放宽标准。

| 后端 / kernel | 耗时 µs | SM吞吐 % | DRAM吞吐 % | L2吞吐 % |
|---|---:|---:|---:|---:|
| FA2 / kernel0 | 1907.808 | 55.363 | 2.102 | 19.837 |
| FA3 / kernel0 | 2.752 | 0.006 | 0.052 | 0.391 |
| FA3 / kernel1 | 918.976 | 79.462 | 4.517 | 22.351 |

FA2为单个attention kernel；FA3第0项为变长参数准备，第1项为attention本体。FA2与FA3均通过FP64抽样及保存参考输出误差门槛；FA3与保存输出逐值一致，FA2未逐值一致但通过预设容差。硬件数据表明此输入下FA3的SM吞吐更高、采集耗时更短；这是同卡独立kernel诊断，仍不能认领整模型加速。

[Attention对照原始指标、数值误差与哈希](/capability-attention-pair-audit.json)

## 算子工作量、访存与VAE布局反证

18个已保存算子的形状、stride、dtype和输入哈希已逐项核对。线性算子按2×M×N×K计算乘加FLOPs；卷积按2×输出元素数×每组输入通道×卷积核体积计算，隐式padding为0，显式填充的输入仍按稠密工作量计数，bias加法单列。转换算子不赋予浮点FLOP率。

这些是算子标称工作量，不能冒充硬件实际执行指令数或剔除零填充后的模型有效工作量。NCU的DRAM流量与耗时统一原生单位后求和；下图左侧是独立回放工作量/流量散点，没有实测带宽上界，因此不是完整Roofline。不同输入、物理GPU和采集回放影响仍需保留。

![独立算子工作量与VAE局部布局对照](/operator-work-traffic.png)

[逐算子数据JSON](/operator-work-traffic.json) · [CSV](/operator-work-traffic.csv) · [矢量图](/operator-work-traffic.svg)

### VAE更细指标与局部对照

代表卷积补采Scheduler、Warp State和Source Counters：寄存器与shared memory理论驻留上限均为1 block，未记录到spill；平均可发射warp约0.199，scheduler发射活跃约18.48%。long scoreboard与barrier等待突出。这支持继续调查延迟隐藏、局部性与同步，不能只凭低DRAM带宽否定访存问题，也不能把warp等待指标当作墙钟占比。

调度记录说明：本次详细采集在正常结束附近约0.27秒内被记录为user unknown，3次观察，无警告、无终止信号，最终为process exited。任务启动身份与输出/计数器审核通过，但不能声称这次运行没有调度异常记录；退出清理竞态只是可能解释。

| 局部实现 | 三批均值 ms | 输出校验 |
|---|---|---|
| 原版 | 0.488351, 0.484531, 0.484496 | 与保存输出逐值一致 |
| channels-last，含全部转换 | 0.595068, 0.594726, 0.594358 | 与保存输出逐值一致 |

候选/原版耗时比为1.2242，即候选约慢22.42%。每次调用都包含输入和权重转布局、卷积及输出转回连续NCDHW；没有省略转换成本。此直接改布局的候选不进入整模型加速认领。

右图数据来自关闭profiler的局部eager CUDA-event计时：每模式预热20次，三个交错批次、每批300次，共900次。是一次进程内的相关样本，不是三次独立启动，也不是整模型尾延迟验收。反例不证明所有布局优化均无效；跨层保持布局或权重预打包需重新定义公平边界并另测。

[详细硬件计数器](/vae-detailed-counters.json) · [详细采集审核与调度备注](/capability-vae-detailed-audit.json) · [局部A/B审核](/vae-layout-local-ab-audit.json)

## 本机连续读写带宽校准

已完成同一物理GPU上三次独立进程启动，每次8种配置，每配置3批×100次未插桩CUDA-event计时，共24个配置/启动点、7200次计时。使用FP32 copy（读A、写Y）与triad（读A/B、写Y），BLOCK为1024或4096、num_warps为4；A=1、B=2，完整输入和每配置最终输出均保存并逐值核验。每次计时循环不保存每次中间输出，避免改变测试负载。

单数组256 MiB或1 GiB，均大于本卡60 MiB L2的4倍；三次进程和独立NCU试采的物理UUID已核对一致。未锁定时钟；遥测已保存。这是所测连续访问kernel达到的带宽参考，不是保证的硬件峰值、其他访问模式的上界或所有GPU的统一校准。

| 操作 | 每数组 MiB | BLOCK | 三次启动带宽 TB/s |
|---|---:|---:|---|
| copy | 256 | 1024 | 4.1136, 4.1136, 4.1151 |
| copy | 256 | 4096 | 3.9135, 3.9135, 3.9153 |
| copy | 1024 | 1024 | 4.2308, 4.2300, 4.2313 |
| copy | 1024 | 4096 | 4.0421, 4.0422, 4.0431 |
| triad | 256 | 1024 | 4.2495, 4.2488, 4.2488 |
| triad | 256 | 4096 | 4.1964, 4.1961, 4.1975 |
| triad | 1024 | 1024 | 4.3496, 4.3498, 4.3504 |
| triad | 1024 | 4096 | 4.2992, 4.2990, 4.2998 |

表中每次启动数值为该启动三批带宽的中位数；带宽按逻辑读写字节数/对应批次事件耗时中位数计算。不同配置的结果约为3.91–4.35 TB/s。窄幅重复结果只描述这三次运行，不证明长期稳定性。

| 1 GiB / BLOCK4096 的NCU核对 | 逻辑字节 GB | DRAM实测 GB | DRAM/逻辑 |
|---|---:|---:|---:|
| copy | 2.147484 | 2.123880 | 0.9890 |
| triad | 3.221225 | 3.200204 | 0.9935 |

独立NCU试采每个kernel只需1次pass，确认大量流量经过DRAM。实测写回字节数与逻辑输出字节数存在小差异，缓存及写回观测窗口可能影响，不能强行视为相等。L2 sector计数保留原生单位，未冒充算法字节数。NCU耗时与无profiler校准分开保存。

前文“尚缺本机带宽参考”现已补上这一连续访问参考；完整Roofline仍需与算子匹配的计算/带宽校准、L2口径以及完整模型有效FLOPs。现有GEMM校准来自另一物理GPU，不把两块卡数据直接当作同卡硬上界。

[三次启动审核与原始文件哈希](/local-bandwidth-audit.json) · [NCU流量核对](/local-bandwidth-ncu-audit.json) · [物理GPU身份审核](/local-bandwidth-gpu-identity.json)

## 同卡计算参考与代表算子位置

三次新的独立启动已完成：BF16、FP32禁用TF32、FP32允许TF32，分别测4096/8192/12288方阵，共27个配置/启动点。每点预热20次、计时100次，之后另采实际kernel trace。完整A/B/输出矩阵均保存；4×4输出使用完整对应行/列做FP64抽样审核，通过既定0.03/0.03 allclose与1%相对L2门槛。该抽样不等于完整矩阵高精度验证。

三次进程与前述带宽校准的物理UUID一致，因此前文“计算与带宽来自不同卡”的限制在这一新增参考组中已解决；旧校准保留为历史数据。参考计算吞吐取每次启动三个方阵中的最大值，再取三次最大值的中位数；仍只是所测kernel达到的参考，不是保证的硬件峰值。

| 配置策略 | 三次启动最大 TFLOP/s | 参考 TFLOP/s |
|---|---|---:|
| BF16 | 787.057, 787.706, 789.081 | 787.706 |
| FP32，禁止TF32 | 51.287, 51.276, 51.250 | 51.276 |
| FP32，允许TF32 | 397.848, 397.443, 397.425 | 397.443 |

![同卡代表算子的实测参考图](/same-gpu-reference.png)

18个NCU算子的原始PCI地址已核对到带时间戳的PCI/UUID表，15个位于校准卡。图中进一步仅纳入7个BF16线性与5个kernel名称明确含TF32的VAE算子；另外3个来自其他卡、2个算术类型尚未充分确认、1个类型转换无定义的乘加FLOPs，均明确排除，不强行归一化。

虚线按min(实测计算参考, 算术强度×实测连续访问带宽参考)构造。点的位置采用NCU累计kernel耗时和DRAM流量，虚线来自关闭profiler的校准；它们用于诊断参照，不能当作同口径加速对比。数值为标称MAC工作量，含已显式填充的输入与多kernel转换成本，非整模型MFU。

观察：部分较大BF16线性算子接近所选参考，而小M=3算子差异明显；5个纳入的VAE算子约为所选参考的4.8%–54.2%。这些差距给出调查方向，并不证明某一stall是唯一原因，更不直接等于可获得的整模型加速。

峰值口径另已核对：[NVIDIA H200官方规格](https://www.nvidia.com/en-us/data-center/h200/)区分SXM/NVL，列出的Tensor Core数字带稀疏性脚注。不能把该稀疏数值直接当成本实验稠密运算的MFU分母，也不凭L20X驱动名称选择峰值。

[同卡计算审核](/same-gpu-gemm-calibration-audit.json) · [计算CSV](/same-gpu-gemm-calibration.csv) · [参考图完整数据与排除原因](/same-gpu-reference.json) · [矢量图](/same-gpu-reference.svg) · [原始PCI映射审核](/operator-ncu-physical-devices.json)

## 新的整模型候选：cross SDPA Flash

选取依据是目标480×832 TP1/SP4配置下6个cross上下文：Flash SDPA的独立回放均通过原定误差门槛，平均耗时约0.04965ms，原版FA3约0.06279ms。native dense在此配置反而更慢，不因其名称而优先选择。已有cuDNN cross整模型候选未通过数值验收，失败证据继续保留。

实际SDPAImpl的6个输入验证现已通过CPU复核，trace包含PyTorch Flash kernel，且固定Flash-only上下文禁止自动回退。候选只替换cross角色的后端配置及对应SDPAImpl CUDA实现，self paged FA路径保持原始源码。

整模型协议：v2参考和候选共用调度器分配的GPU2/3/4/6；原GPU4/5/6/7队列因资源等待在模型启动前结束，旧记录保留。原版固定源码、seed42、eager模式，每侧两轮；先运行新参考并与已有可重复参考比较40份完整输出，要求逐值一致，再执行候选。最终候选沿用latent与归一化像素原有门槛，不放宽阈值。

当前流程状态：comparison_completed。排队或运行状态不是验证成功。

[六个实际实现验证的审核](/cross-sdpa-flash-impl-audit.json)

新GPU组参考与历史eager参考：40份完整输出，逐值一致40份；数值门槛通过。此项只验证参考输出，不代表候选已通过或可将不同GPU组耗时合并。

新整模型比较：40份完整输出，门槛通过0份，逐值一致0份。候选未通过原定数值门槛，不认领有效加速。

来源与单次计时审核：同卡新参考 11.7731 FPS，候选 11.7881 FPS；平均CPU-ready交付间隔分别 1019.273 / 1017.975 ms。每侧只有4个稳态间隔，且候选数值未通过，不能解释为有效加速。

日志名SDPA对应配置TORCH_SDPA：注册表指向SDPABackend，get_name返回SDPA。已按源码修正审核器名称匹配；原失败审核保留。四个worker选择及Flash-only实现安装可追溯，但本次无整模型原生kernel trace。

[完整输出误差](/cross_sdpa_model_comparison_v2.json) · [来源与计时审核](/cross_sdpa_model_audit_v2.json) · [日志名称修正依据](/cross_sdpa_log_name_correction_v2.json)

语义排查：预热和测量轮的latent、pixels均从第0块就未通过门槛；不能把失败只归因于长序列后期漂移。cross调用源码未显式传入掩码或packed metadata，两侧沿用相同缓存K/V、scale与非causal设置。但缺少本候选逐层同输入对照，尚不能判定是数值放大还是内部实现差异。六个局部probe使用裁剪后的连续张量，不等于完整运行时stride与全层覆盖。

[源码快照索引及逐块误差](/cross_sdpa_semantic_review_v2.json)

后续真实输入诊断：completed_pending_audit，已保存160/160次匹配调用。rank0、B1/H10/K512、非causal筛选；每次使用同一实时Q/K/V计算原版与SDPA Flash输出，并保存FP64抽样参考。原版输出继续用于模型。该诊断有额外计算与保存开销，全部耗时排除于性能结论，尚需完成独立审核。

真实输入诊断已完成独立CPU复算：160次局部候选/原版比较全部通过原门槛，最大相对L2误差0.001090；两侧对FP64抽样参考也全部通过。保存约6.19GB张量。诊断运行最终40份输出与新参考逐值一致，说明本次影子计算未改变这些保存输出。

这缩小了排查范围：在原版轨迹首块的160次已采集调用中，局部差异较小；但候选整模型轨迹会随每次替换而变化，本诊断没有重现该轨迹，也未覆盖所有rank和后续块。因此仍不能证明累积舍入是唯一根因，不能接受候选加速。

[160次同输入独立审核](/cross_sdpa_shadow_audit_v1.json) · [诊断输出一致性](/cross_sdpa_shadow_model_invariance_v1.json)

候选自身轨迹诊断：completed_pending_audit，已保存160/160次调用。全部rank继续使用SDPA Flash候选输出，仅rank0前160次目标调用额外计算原版和FP64抽样。保存完整输入、两种输出；待独立审核和轨迹对应，不将此运行耗时纳入性能结论。

候选轨迹诊断已通过独立复算，最终40份输出与未插桩候选逐值一致。其160次同输入原版/候选比较和两侧FP64抽样仍全部通过，局部最大相对L2约0.001089。两条轨迹按调用顺序配对，形状、stride、dtype与缩放均一致，K/V在全部160对中逐值相同；Q第一对相同，从第2次调用出现差异，最大相对L2约0.060016。差异并非单调增长。

这支持“局部小差异伴随模型状态分化”的解释，但不证明舍入是唯一原因：尚缺层名、全部rank及其他操作的逐层受控归因；跨轨迹Q差异也不是同输入kernel误差。整模型数值失败结论保持。

![同输入局部误差与跨轨迹Q差异](/cross_shadow_trajectories.png)

[候选轨迹独立审核](/cross_sdpa_candidate_shadow_audit_v1.json) · [候选输出不变检查](/cross_sdpa_candidate_shadow_model_invariance_v1.json) · [完整配对数据](/cross_shadow_trajectory_comparison_v1.json) · [矢量图](/cross_shadow_trajectories.svg)

## 小工作集带宽：尚不能作为L2峰值

GPU6三次独立进程测试2/8/16MiB数组、copy/triad与两种block配置，共36个配置启动点、10800个事件样本；完整输入和最终输出、PTX/cubin、逐次计时通过审核。三个16MiB数组总容量小于报告的60MiB L2，但容量关系并不证明所有访问命中缓存。

另用NCU在16MiB数组上关闭缓存清空，分别采集copy/triad。copy为7.584μs，DRAM读写合计7.566MB；triad为12.608μs，DRAM读写合计40.725MB，对应算法字节33.554/50.332MB。仍有明显DRAM流量；写回统计窗口和缓存状态可能影响计数，不能将算法字节直接当成L2物理流量。

未插桩逐次提交事件计时与NCU单kernel时长属于不同测量条件。小kernel事件区间可能包含提交间隔，尚不能定量归因为CPU开销，也不能把本组算法带宽作为L2上限。后续需要设备端连续执行与流量验证，当前L2参考缺口保留。

[36点与10800计时审核](/warm_working_set_audit_v1.json) · [原生NCU指标和单位](/warm_working_set_ncu_audit_v1.json)

## 小工作集提交方式对照

同一进程、同一输入和kernel，三批交替普通提交与包含100次调用的CUDA Graph，共三次独立启动、72个配置/方法/启动点、21600个事件区间。Graph原始总时长完整保留；下表Graph值为总时长除以100后的平均值，各列再取三启动中位数。完整输入及每配置最终输出通过审核；未保存每次调用或每种方法的独立输出快照。

| 模式 | 数组MiB | block | 普通提交μs | Graph平均μs |
|---|---:|---:|---:|---:|
| copy | 2 | 1024 | 18.528 | 1.539 |
| copy | 2 | 4096 | 19.728 | 1.628 |
| copy | 8 | 1024 | 19.712 | 3.126 |
| copy | 8 | 4096 | 19.232 | 3.226 |
| copy | 16 | 1024 | 19.920 | 5.269 |
| copy | 16 | 4096 | 19.776 | 5.771 |
| triad | 2 | 1024 | 18.272 | 1.693 |
| triad | 2 | 4096 | 20.464 | 1.927 |
| triad | 8 | 1024 | 19.968 | 3.853 |
| triad | 8 | 4096 | 19.008 | 4.193 |
| triad | 16 | 1024 | 23.232 | 11.537 |
| triad | 16 | 4096 | 23.328 | 10.884 |

提交方式显著改变小工作集的测量结果。Graph还可能改变缓存状态和执行间隔，不能把差额全部归为CPU开销；Graph平均值不是独立kernel尾延迟。算法带宽最高约6.53TB/s，仍不是纯L2峰值，未据此填补完整L2 Roofline。该微基准结果不代表整模型加速。

[三启动配对数据、设备UUID及校验](/warm_graph_pair_audit_v1.json)

## 已记录cross-attention的矩阵运算量

按实际记录的B/Q/K/H/D与调用次数，计算QKᵀ和PV两次矩阵乘：每调用4×B×Q×K×H×D，乘和加各算1 FLOP。下表为诊断记录范围内各rank求和后的总运算量，不是每秒吞吐；不同采集的轮次/调用数不一定相同，不能直接排名。

| 采集 | cross调用数 | 层前缀数 | 标称矩阵运算量（万亿FLOP） |
|---|---:|---:|---:|
| capture_cross_480x832_tp1_sp4_v1 | 16000 | 40 | 196.293 |
| capture_dit_only_704x1280_tp1_sp1_v2 | 2000 | 40 | 221.459 |
| capture_evening_480x832_tp1_sp1_v1 | 4000 | 40 | 196.293 |
| capture_evening_480x832_tp2_sp1_v1 | 8000 | 40 | 196.293 |
| capture_evening_480x832_tp2_sp2_v1 | 16000 | 40 | 196.293 |
| capture_evening_480x832_tp4_sp1_v1 | 16000 | 40 | 196.293 |
| capture_evening_704x1280_tp1_sp2_v1 | 8000 | 40 | 442.919 |
| capture_evening_704x1280_tp2_sp1_v1 | 8000 | 40 | 442.919 |
| capture_evening_704x1280_tp2_sp2_v1 | 16000 | 40 | 442.919 |
| capture_evening_704x1280_tp4_sp1_v1 | 16000 | 40 | 442.919 |
| capture_retry_480x832_tp1_sp2_v1 | 8000 | 40 | 196.293 |
| capture_retry_704x1280_tp1_sp4_v1 | 16000 | 40 | 442.919 |

480×832 TP1/SP4代表采集已从原始JSONL独立重数：16000次cross调用，共196.293万亿矩阵FLOP，其中denoise157.035、commit39.259。数字包括记录中的重复轮次，不能除以另一运行的耗时来计算MFU。

本项不含softmax、缩放、投影、归一化、self-attention或VAE。self逐调用cu_seqlens实际数值缺失，不能用分配的K长度或max_seqlen_k替代有效长度；完整模型FLOPs/MFU缺口保留。

[12个采集的运算量与形状文件SHA](/cross_attention_workload.json) · [代表采集原始行复算](/cross_attention_workload_raw_recheck.json)

## 逐调用有效Attention长度补充

480×832、TP1/SP4、eager，两轮各10个chunk：四个rank的40层均覆盖，共32000次调用（self与cross各16000）。保存实际cu_seqlens数值、形状、步数/阶段/层上下文及逐调用运算量账本；40份最终模型输出与独立参考逐值一致。

两轮全部记录的self QKᵀ/PV矩阵运算量为8074.101万亿FLOP，cross为196.293万亿FLOP。若错误使用分配长度，self会算成11363.549万亿FLOP，高估40.74%。这证明本配置必须使用有效长度统计，不能把缓存容量作为实际工作量。

本次同步读取长度会影响计时，不能据此计算无干扰吞吐或MFU。运算量仍不含softmax、投影、归一化和VAE；仅补齐此480×832配置，两轮统计总量不能作为单帧工作量。中间仅保留元数据，未保存每次完整QKV；最终40份输出已保存。

[长度与覆盖审核](/attention_length_audit_v1.json) · [40份输出一致性](/attention_length_model_invariance_v1.json) · [逐调用账本（gzip）](/attention_length_call_work_v1.jsonl.gz)

## 长尾区间对会话相关性的敏感性

复用六次既有运行，每次250个测量会话、每会话4个CPU-ready间隔。原报告已经按整个会话重采样；这次进一步采用循环连续会话块，块长1、5、10、25，每种2000次重采样，始终保留同会话的4个间隔。原始重采样索引与结果全部保存，未新增GPU运行。

下表是P99估计的2.5%–97.5%重采样分位区间。连续块方法仍依赖足够平稳的序列假设，不能由它本身证明平稳性或未来95%覆盖率；不同块长全部保留，不选择最窄区间作为结论。

| 运行（编号从0开始） | 会话均值lag-1相关 | 单会话块 P99区间 ms | 25会话块 P99区间 ms |
|---|---:|---|---|
| 480x832_r0 | 0.168 | 836.270–842.137 | 836.271–841.136 |
| 480x832_r1 | 0.393 | 854.247–886.014 | 853.724–885.880 |
| 480x832_r2 | 0.297 | 845.619–861.729 | 845.278–861.729 |
| 704x1280_r0 | 0.451 | 2496.570–2703.041 | 2490.742–2703.041 |
| 704x1280_r1 | 0.295 | 2513.651–2753.535 | 2479.617–2792.246 |
| 704x1280_r2 | 0.130 | 2471.549–2481.619 | 2471.060–2478.003 |

![长尾区间对连续会话块长的敏感性](/tail-block-sensitivity.png)

704×1280第2次运行（r1）的P99区间由约2514–2754ms扩展到2480–2792ms，其他运行不一定变宽。相邻会话均值存在不同程度的相关性；部分运行前50与后50会话均值也不同。这些是描述性证据，不能据此归因于某个硬件因素。

因此，1000个间隔满足样本数量门槛，却不自动建立独立性、长期稳定性或服务SLA。该数据仍来自固定输入重复请求，且止于CPU-ready；不能替代多输入/多seed、完整输出数值验收和实际客户端交付测试。

[全部4种块长的数据与哈希](/tail-block-sensitivity.json) · [独立核对记录](/tail-block-sensitivity-audit.json) · [矢量图](/tail-block-sensitivity.svg)

## 同次运行的主要矩阵运算量账本

480×832 TP1/SP4 eager两轮诊断完成。所有记录按物理GPU映射核对；四rank各40800次线性调用、各1996次VAE卷积，self/cross各16000次，另有20次GPU输出类型转换。40份最终输出与独立参考逐值一致。

| 类别 | 已记录调用数 | 两轮标称矩阵运算量（万亿FLOP） |
|---|---:|---:|
| cross-attention核心 | 16000 | 196.293 |
| self-attention核心 | 16000 | 8074.101 |
| DiT线性层 | 163200 | 15344.762 |
| 输出类型转换 | 20 | 未分配算术FLOPs |
| VAE卷积 | 7984 | 724.657 |

上述矩阵工作量来自同一次运行。DiT记录仅含linear，避免与Attention核心矩阵乘重复计数。卷积按实际输出及权重形状计算标称密集运算量，包含空间分片/填充可能引入的工作；不等同于硬件实际执行指令数。

仍未包含偏置、激活、归一化、softmax、位置变换等算术，也不涵盖编码和网络。输出转换未分配FLOPs，不表示没有成本。同步采集和逐调用日志会改变耗时，不以本次诊断时长计算性能或MFU。中间仅保存元数据，最终40份输出完整保存。

[同次工作量与物理GPU映射](/recorded_model_mac_work_v1.json) · [算子记录审核](/operator_metadata_audit_v1.json) · [Attention长度审核](/operator_metadata_attention_audit_v1.json) · [最终输出一致性](/operator_metadata_model_invariance_v1.json)

## 代表NCU输入与逐调用结构的对应范围

将既有18个NCU输入快照，与新诊断中每次算子的嵌套参数、输入/输出形状、步长、dtype及标量参数逐项对照。跨rank比较时仅去掉device字段。

| 类别 | 实际不同签名数 | 已有同签名NCU输入数 | 同签名调用数 / 已记录调用数 |
|---|---:|---:|---:|
| dit | 12 | 8 | 114800 / 163200 |
| output | 2 | 2 | 20 / 20 |
| vae | 17 | 8 | 2872 / 7984 |

仍缺4种线性层和9种卷积签名的代表硬件输入；线性缺口包括序列分片后的1170长度投影/MLP，以及最终64通道投影。卷积缺口包含较大空间分辨率和后续时间块。应据此补采真实输入，而不是把最初8种形状视作全覆盖。

结构一致不证明输入值相同、同一调用或同一native kernel，也不能把旧回放计时乘调用次数作为整模型耗时。表中是结构对应数量，不是硬件计数器采集调用覆盖率或FLOP加权覆盖率。

[全部签名、频次、匹配输入及源文件SHA](/operator_replay_signature_coverage_v1.json)

## 结构对应的标称工作量权重

将上述签名按逐调用密集MAC工作量加权，并与独立算子账本总量复核：旧代表输入对应线性层工作量约59.06%，VAE卷积约5.09%。其余约40.94%与94.91%的标称工作量所在签名，尚无旧NCU输入对应。后续13种输入补采正针对这个缺口。

这个比例仅说明已记录矩阵工作量分布，不是硬件采集覆盖率、时间占比或真实指令数。Attention核心和非矩阵算术不在分母内；不能由VAE缺口大直接推断VAE占整模型耗时大。

[逐签名工作量、排序与复核](/signature_work_coverage_v1.json)

## 新增13种真实输入的硬件结果

补采4种线性层和9种卷积输入共1.687GB，40份最终模型输出与参考逐值一致。随后通过普通用户gpu run在GPU6完成13项隔离NCU回放，共48个kernel；13份保存的回放输出经CPU独立比对均逐值一致，权限/CVD、计数器来源和各采集PID的调度违规历史检查通过。

| 类别 | 输入形状 | 输出形状 | kernel数 | NCU累计μs | 标称矩阵TFLOP/s |
|---|---|---|---:|---:|---:|
| dit | [1, 1170, 5120] | [1, 1170, 13824] | 1 | 215.520 | 768.48 |
| dit | [1, 1170, 13824] | [1, 1170, 5120] | 1 | 220.608 | 750.76 |
| dit | [1, 1170, 5120] | [1, 1170, 5120] | 1 | 85.696 | 715.81 |
| vae | [1, 192, 6, 242, 106] | [1, 192, 4, 240, 104] | 5 | 872.416 | 227.81 |
| vae | [1, 96, 6, 482, 210] | [1, 96, 4, 480, 208] | 5 | 1672.160 | 118.86 |
| vae | [1, 384, 4, 122, 54] | [1, 384, 2, 120, 52] | 5 | 486.240 | 204.37 |
| vae | [1, 192, 4, 122, 54] | [1, 384, 2, 120, 52] | 5 | 253.024 | 196.37 |
| vae | [1, 384, 4, 120, 52] | [1, 768, 2, 120, 52] | 6 | 168.224 | 131.27 |
| vae | [1, 96, 6, 482, 210] | [1, 3, 4, 480, 208] | 5 | 704.128 | 8.82 |
| dit | [1, 4680, 5120] | [1, 4680, 64] | 1 | 17.600 | 174.27 |
| vae | [1, 384, 3, 60, 26] | [1, 768, 1, 60, 26] | 5 | 51.712 | 53.38 |
| vae | [1, 192, 2, 120, 52] | [1, 384, 2, 120, 52] | 3 | 43.936 | 41.88 |
| vae | [1, 96, 3, 482, 210] | [1, 3, 1, 480, 208] | 5 | 284.000 | 5.47 |

例如192通道大卷积：布局转换68.352+5.216+39.008μs、卷积706.240μs、后续逐元素53.600μs，总计872.416μs。卷积kernel的SM吞吐约78.23%，两次主要布局转换L2吞吐约82.80%和85.50%。计算与布局成本都存在，不能仅看卷积kernel推断整个算子；此前局部channels-last负结果也不能直接推广到这些新增形状。

表中速度以标称密集矩阵运算量除以NCU累计kernel时长，包含本次布局/逐元素kernel，不是无插桩延迟或模型MFU。即使输入签名与逐调用记录匹配，仍不能把隔离回放时间乘调用数作为整模型墙钟耗时。

[13项硬件指标与原始报告校验](/missing_operator_ncu_audit_v1.json) · [形状/工作量/DRAM流量](/missing_operator_work_traffic_v1.json) · [真实快照审核](/missing_operator_capture_audit_v1.json) · [模型输出一致性](/missing_operator_model_invariance_v1.json) · [补齐后结构对应](/operator_replay_signature_coverage_v2.json)

## 新增卷积的计算与转换拆分

按明确的native kernel名称分组，并核对每项分组时长之和与原始NCU总量一致。比例仅为隔离回放累计kernel时间占比，不是模型墙钟占比或可恢复加速。

| 输入形状 | 输出形状 | 总μs | 计算% | 布局转换% | 其他% |
|---|---|---:|---:|---:|---:|
| [1, 192, 6, 242, 106] | [1, 192, 4, 240, 104] | 872.416 | 81.0 | 12.9 | 6.1 |
| [1, 96, 6, 482, 210] | [1, 96, 4, 480, 208] | 1672.160 | 80.1 | 13.4 | 6.4 |
| [1, 384, 4, 122, 54] | [1, 384, 2, 120, 52] | 486.240 | 87.9 | 9.0 | 3.1 |
| [1, 192, 4, 122, 54] | [1, 384, 2, 120, 52] | 253.024 | 82.7 | 11.6 | 5.8 |
| [1, 384, 4, 120, 52] | [1, 768, 2, 120, 52] | 168.224 | 48.5 | 33.0 | 18.6 |
| [1, 96, 6, 482, 210] | [1, 3, 4, 480, 208] | 704.128 | 76.2 | 22.9 | 0.8 |
| [1, 384, 3, 60, 26] | [1, 768, 1, 60, 26] | 51.712 | 36.1 | 51.9 | 11.9 |
| [1, 192, 2, 120, 52] | [1, 384, 2, 120, 52] | 43.936 | 45.5 | 0.0 | 54.5 |
| [1, 96, 3, 482, 210] | [1, 3, 1, 480, 208] | 284.000 | 70.6 | 28.2 | 1.3 |

较大卷积的计算比例约80%–88%；输出[1,768,1,60,26]的时间维卷积则约51.9%是布局转换。该差异提示优化需要按真实形状验证，不能将单一布局结论推广到全部VAE。

[完整kernel分类及时间复核](/missing_vae_components_v1.json)

## 时间维卷积布局A/B：仍未获得收益

选择新增NCU中布局转换比例最高（约51.9%）的时间维卷积，使用保存的真实输入做无profiler局部A/B。保持相同NCDHW输入输出边界，候选每次都计入输入/权重channels-last打包和输出连续化。

单次进程内每种20次预热，交替三批各300次CUDA事件计时；两种完整输出与原始参考逐值一致。原方案三批均值约51.39、54.97、51.33μs，候选60.58、60.17、60.31μs。平均52.56对60.35μs，候选慢约14.82%，不接受为优化收益。

这并不否定跨多层保持布局或提前打包的可能性，但那些改变了本次边界，仍需独立数值与全成本验证。本次只有一个输入、一个进程，三批不是三次独立启动，也不是模型尾延迟测试。

[完整原始计时/输出审核](/vae_layout_temporal_ab_audit_v1.json)

## KV增长期间的逐chunk矩阵工作量

复用同次480×832 TP1/SP4逐调用数据，按记录的start_frame分组，包含每chunk四次denoise和一次commit，四rank只各计一次；两轮总量除以2，得到单轮平均。逐组结果已与原始工作量账本总量核对。

| 记录start_frame | 有效K | 线性层 万亿FLOP | self 万亿FLOP | cross 万亿FLOP | self占已统计DiT矩阵工作量 |
|---:|---:|---:|---:|---:|---:|
| 0 | 4680 | 767.24 | 89.71 | 9.81 | 10.35% |
| 3 | 9360 | 767.24 | 179.42 | 9.81 | 18.76% |
| 6 | 14040 | 767.24 | 269.14 | 9.81 | 25.73% |
| 9 | 18720 | 767.24 | 358.85 | 9.81 | 31.59% |
| 12 | 23400 | 767.24 | 448.56 | 9.81 | 36.60% |
| 15 | 28080 | 767.24 | 538.27 | 9.81 | 40.92% |
| 18 | 28080 | 767.24 | 538.27 | 9.81 | 40.92% |
| 21 | 28080 | 767.24 | 538.27 | 9.81 | 40.92% |
| 24 | 28080 | 767.24 | 538.27 | 9.81 | 40.92% |
| 27 | 28080 | 767.24 | 538.27 | 9.81 | 40.92% |

实际有效K在start_frame=15达到28080，之后此记录范围内保持不变。self矩阵工作量增长至起始的6倍，线性层及cross矩阵工作量保持固定；因此已统计矩阵工作量中的self占比由10.35%升至40.92%后稳定。

分母仅含已记录线性层与Attention核心矩阵乘，不含归一化、激活、softmax、位置变换、VAE等。运算量占比不能直接解释墙钟延迟、通信重叠或MFU；该规律只在本配置和已记录序列范围成立。

[逐chunk数据、调用数与源文件校验](/dit_work_by_chunk_v1.json)

## 704×1280同口径诊断与分辨率对照

新无插桩eager参考和元数据诊断在同批GPU0/1/2/3完成；40份输出逐值一致，像素704×1280、latent88×160，来源/输入/分配/物理映射审核通过。两种分辨率均为两轮、TP1/SP4；各有163200次线性、7984次卷积、32000次Attention及20次转换。

| 类别 | 704标称矩阵工作量 / 480 |
|---|---:|
| attention_cross | 2.256410 |
| attention_self | 5.091387 |
| dit | 2.256374 |
| vae | 2.256410 |

704有效K由10560增至63360后保持；self在已统计DiT矩阵工作量中的比例由20.67%升至60.98%。480对应10.35%至40.92%。线性层、cross和卷积约随像素面积增至2.256倍；self核心矩阵工作量约增至5.091倍。这里是实际长度/形状推导的标称矩阵工作量，不是实测耗时比或完整MFU。

结构对照发现：旧31种NCU代表输入仅匹配704中的3/12种线性签名（1200次调用），0/17种卷积及0/2种转换；仍缺28种高分辨率签名的代表硬件输入。不能把480已补齐的结构覆盖外推到704。

首次来源审核误将参考脚本成功状态pilot_passed要求为completed，已按原脚本定义修正并重跑CPU审核；失败日志保留，没有重跑GPU或放宽数值门槛。

[分辨率工作量对照](/resolution_recorded_work_comparison_v1.json) · [704逐chunk工作量](/dit_work_by_chunk_704_v1.json) · [704结构覆盖](/operator_replay_signature_coverage_704_v1.json) · [来源/尺寸/输出审核](/operator_metadata_704_provenance_audit_v1.json) · [全部模型输出比较](/operator_metadata_704_model_invariance_v1.json) · [同次工作量与物理映射](/recorded_model_mac_work_704_v1.json)

## 704新增28项硬件回放：含一项调度备注

高分辨率28个真实快照共5.302GB，40份模型输出与新参考逐值一致。普通用户gpu run分配GPU0后完成28项独立NCU回放，共90个kernel；保存的28份回放输出均经CPU独立比较逐值一致，计数器来源、UID和CVD核对通过。

严格调度历史审核发现PID2019660一条unknown/foreign记录：首次05:20:53.219 UTC，末次05:20:54.101 UTC，3次观察，未发警告/信号，随后记录为process exited。该项采集结束时间05:20:53.252 UTC，身份记录UID2005/CVD0正确；异常发生在结束附近，但不能据此证明是退出竞态。原严格审核失败日志保留，后续只将计数器/输出检查与调度归属异常分别记录，不宣称调度无异常。其余27项没有匹配的自身PID违规记录。

| 类别 | 输入形状 | 输出形状 | kernel数 | NCU累计μs | 标称矩阵TFLOP/s | 调度备注 |
|---|---|---|---:|---:|---:|---|
| dit | [1, 2640, 5120] | [1, 2640, 5120] | 1 | 185.344 | 746.78 | 无自身PID记录 |
| dit | [1, 2640, 5120] | [1, 2640, 15360] | 1 | 518.112 | 801.44 | 无自身PID记录 |
| dit | [1, 2640, 13824] | [1, 2640, 5120] | 1 | 472.896 | 790.26 | 无自身PID记录 |
| dit | [1, 2640, 5120] | [1, 2640, 13824] | 1 | 451.040 | 828.56 | 无自身PID记录 |
| dit | [1, 2640, 5120] | [1, 2640, 5120] | 1 | 183.520 | 754.21 | 无自身PID记录 |
| vae | [1, 96, 6, 706, 322] | [1, 96, 4, 704, 320] | 5 | 3677.952 | 121.93 | 无自身PID记录 |
| vae | [1, 192, 6, 354, 162] | [1, 192, 4, 352, 160] | 5 | 2130.848 | 210.46 | 无自身PID记录 |
| dit | [1, 10560, 5120] | [1, 10560, 5120] | 1 | 650.336 | 851.33 | 无自身PID记录 |
| vae | [1, 384, 4, 178, 82] | [1, 384, 2, 176, 80] | 5 | 1129.536 | 198.51 | 无自身PID记录 |
| vae | [1, 384, 3, 90, 42] | [1, 384, 1, 88, 40] | 5 | 180.320 | 155.44 | 无自身PID记录 |
| dit | [1, 10560, 1536] | [1, 10560, 5120] | 1 | 208.928 | 794.98 | 结束附近unknown记录 |
| vae | [1, 192, 4, 178, 82] | [1, 384, 2, 176, 80] | 5 | 603.968 | 185.63 | 无自身PID记录 |
| vae | [1, 384, 4, 176, 80] | [1, 768, 2, 176, 80] | 5 | 299.136 | 166.57 | 无自身PID记录 |
| dit | [10560, 144] | [10560, 5120] | 1 | 53.472 | 291.21 | 无自身PID记录 |
| vae | [1, 96, 3, 706, 322] | [1, 96, 1, 704, 320] | 5 | 1006.016 | 111.44 | 无自身PID记录 |
| vae | [1, 192, 3, 354, 162] | [1, 192, 1, 352, 160] | 5 | 562.496 | 199.31 | 无自身PID记录 |
| vae | [1, 384, 3, 178, 82] | [1, 384, 1, 176, 80] | 6 | 585.408 | 191.51 | 无自身PID记录 |
| vae | [1, 96, 6, 706, 322] | [1, 3, 4, 704, 320] | 5 | 1515.040 | 9.25 | 无自身PID记录 |
| dit | [1, 10560, 5120] | [1, 10560, 64] | 1 | 31.136 | 222.27 | 无自身PID记录 |
| vae | [1, 384, 3, 88, 40] | [1, 768, 1, 88, 40] | 6 | 74.752 | 83.32 | 无自身PID记录 |
| vae | [1, 192, 2, 176, 80] | [1, 384, 2, 176, 80] | 3 | 80.960 | 51.29 | 无自身PID记录 |
| vae | [1, 192, 3, 178, 82] | [1, 384, 1, 176, 80] | 6 | 306.400 | 182.95 | 无自身PID记录 |
| vae | [1, 16, 3, 90, 42] | [1, 384, 1, 88, 40] | 4 | 42.368 | 27.56 | 无自身PID记录 |
| vae | [1, 96, 3, 706, 322] | [1, 3, 1, 704, 320] | 5 | 502.560 | 6.97 | 无自身PID记录 |
| vae | [1, 192, 1, 176, 80] | [1, 384, 1, 176, 80] | 2 | 36.704 | 56.57 | 无自身PID记录 |
| vae | [1, 16, 3, 88, 160] | [1, 16, 3, 88, 160] | 2 | 10.496 | 2.06 | 无自身PID记录 |
| output | [1, 3, 9, 704, 1280] | [1, 3, 9, 704, 1280] | 1 | 64.640 | 未定义 | 无自身PID记录 |
| output | [1, 3, 12, 704, 1280] | [1, 3, 12, 704, 1280] | 1 | 84.224 | 未定义 | 无自身PID记录 |

结合旧输入，704此逐调用记录的12种线性、17种卷积和2种转换结构均已找到代表NCU输入；两个分辨率合计59种代表输入、183个kernel记录，包含上述一项调度备注。仍不证明所有调用/张量值的kernel等价或整模型MFU，隔离NCU计时不可乘调用数当模型墙钟时间。

[计数器/输出审核与调度备注](/missing_operator_704_ncu_audit_v1.json) · [工作量/流量和逐项备注](/missing_operator_704_work_traffic_v1.json) · [补齐后的704结构覆盖](/operator_replay_signature_coverage_704_v2.json) · [真实输入审核](/missing_operator_704_capture_audit_v1.json) · [模型输出一致性](/missing_operator_704_model_invariance_v1.json)

## 高分辨率VAE的不同成本构成

以下17个真实卷积输入按native kernel名称分成互斥类别，分组时长之和已与原NCU累计时间核对。左侧是隔离回放累计kernel时间占比，右侧是累计时间的对数坐标；都不表示模型墙钟占比或可恢复加速。

![704×1280 VAE各真实形状的计算与转换构成](/vae_components_704.png)

较大卷积的计算占比约80%–90%。输出[1,768,1,88,40]的时间维卷积则有约44.3%布局转换、12.9%逐元素加法和5.9%工作区初始化；计算仅约36.9%。16→16通道小卷积总计10.496μs，逐元素加法约44.5%；192→384通道另一小算子约44.2%为逐元素加法。高占比也可能对应较小绝对时间，不能据此直接优先做整模型优化。

图表说明需要按真实形状选择优化：大卷积的计算效率、时间维卷积的转换、小算子的后处理，属于不同问题。已完成的两个全成本布局A/B均为负结果，不能把图中的转换比例直接当作channels-last收益；跨层保持布局、权重预打包或融合后处理仍需单独验证。

[完整kernel分组与源证据](/vae_components_704_v1.json) · [矢量图](/vae_components_704.svg)
## 短 KV、长 KV 与 cross-attention 的 FA2/FA3 硬件对照

四组真实输入、八次采集均通过既定数值检查，原始 NCU 报告、完整输出和 FP64 抽样参考已保存。每组在同一 GPU 上依次运行 FA2、FA3。

| 输入 | 有效 KV 长度 | FA2 内核合计 μs | FA3 内核合计 μs | FA2/FA3 耗时比 |
|---|---:|---:|---:|---:|
| 480_short | 4680 | 332.480 | 179.744 | 1.850 |
| 704_short | 10560 | 1594.784 | 821.184 | 1.942 |
| 704_long | 63360 | 9490.592 | 4683.872 | 2.026 |
| 704_cross | 512 | 92.928 | 80.416 | 1.156 |

FA3 合计包括准备 kernel。self-attention 的局部差距大于这组 cross-attention；这是 profiler 下的独立回放结果，不能直接推出整模型加速比例，也不是未插桩延迟测试。短、长 KV 来自不同生成阶段，不能把差异完全归因于长度。

数值检查使用有效 KV，排除分配缓冲区的无效尾部。8 次 FP64 抽样检查及完整输出容差检查均通过；FA3 的三组 self 输出与原捕获值完全一致，但 cross 输出并非逐位一致。局部容差通过不代表整模型输出通过。

证据：[汇总数据](attention_hardware_extension_summary_v1.json)；原始目录：`raw/attention_hardware_extension_v1/`。各输入、输出及报告 SHA256 见 `analysis/attention_hardware_*_audit_v1.json`。

## 同一 Attention 输入的两种计时口径

按真实输入 SHA256 匹配已有未插桩 CUDA event 回放与新增 NCU。它们来自不同运行和 GPU 分配，因此不据此估算 profiler 开销。两类测量的方向一致，但比例并不相同。

| 输入 | 未插桩 FA2/FA3 耗时比 | NCU FA2/FA3 耗时比 |
|---|---:|---:|
| 480_short | 1.822 | 1.850 |
| 704_short | 1.806 | 1.942 |
| 704_long | 1.833 | 2.026 |
| 704_cross | 1.194 | 1.156 |

未插桩结果取原有同进程三批均值，不是三次独立启动。NCU合计包括FA3准备kernel。收益排序仅适用于这些输入，不能替代整模型验收。

高分辨率长KV主kernel的SM吞吐指标由FA2约55.83%升到FA3约77.72%，occupancy由11.92%升到18.75%；DRAM吞吐指标分别约1.22%、2.53%。这些观察支持继续检查执行效率，但低DRAM百分比本身不足以证明计算瓶颈，SM综合指标也不是Tensor Core利用率。

汇总另外列出QK与AV的标称运算量（4QKHD）和实测DRAM流量。该运算量不含softmax等辅助操作，不能当作完整执行FLOPs或模型MFU。证据：[输入匹配及计数器数据](attention_counter_context_v1.json)。

## 输出生产者与消费者边界核对

冻结Case D的rank0对FP32 BCTHW解码输出执行加一、缩放、round、clamp和uint8转换，再同步复制到CPU。转换没有BCTHW到BTHWC的置换；已有uint8传输路径不能重新认领为优化收益。

CPU-ready时间戳记录在`.cpu()`后；随后才检查finite、按配置保存输出、构造`cpu.numpy().tobytes()`并计算SHA256。父进程收到的是chunk信息、时间戳、形状、哈希和内存统计，回复字典未包含像素张量。因此回复时间不是像素IPC交付时间，更不是编码完成或客户端可播放时间。

已为GPU解码输出、GPU量化、D2H、哈希准备、元数据IPC和编码/交付建立六阶段表。张量逻辑字节数按9帧首块或12帧后续块列出，不冒充实测PCIe流量、CPU分配次数或copy engine计数器。编码与真实像素IPC在该冻结基线未执行，服务边界仍需独立实验。

证据：[输出边界与逻辑尺寸清单](output_boundary_inventory_v1.json)。输出融合补采和局部A/B状态见实时证据页，未完成前不报告收益。

## 新增CPU无损媒体往返验证

使用已通过参考核验的480×832连续117帧，按原帧序从BCTHW转为RGB24，再以FFV1 level3、bgr0编码至Matroska。固定4个CPU及4编解码线程，24fps仅为播放时间基准。三个独立子进程编码耗时分别为1.517、1.496、1.665秒，解码为2.417、2.343、2.325秒。

三份视频均经ffprobe确认117帧、832×480、FFV1和24fps。独立重读核对351个解码帧，全部与对应原始RGB字节一致；未发现帧序变化。每份媒体53,725,551字节，未压缩RGB为140,175,360字节。容器文件哈希不同不代表像素不同。

该新增实验使用文件输入输出，耗时包括进程启动和文件I/O，不包括张量准备、模型解码、像素IPC或网络。它是固定无损编码路径的完整性证据，不是浏览器兼容性、有损质量、线上尾延迟或服务吞吐验收。原始RGB、三份视频、三份完整解码RGB、逐帧哈希、命令和日志均已保留。

证据：[媒体往返审核](lossless_media_roundtrip_audit_v1.json)；原始目录：`raw/lossless_media_roundtrip_480_v1/`。

## 输出转换融合局部对照与交接终态

修正版补采已完成：20份真实FP32解码输出全部保留，40份整模型保存输出与参考逐值一致。Triton融合转换对20份真实输入及775个取整边界值均通过独立逐字节检查；CPU参考和原模型uint8输出一致。

| 输入 | GPU-only候选/原版 | 转换+D2H墙钟候选/原版 |
|---|---:|---:|
| 9帧480×832 | 0.292 | 0.762 |
| 12帧480×832 | 0.256 | 0.682 |

包含D2H的局部耗时分别降低23.76%和31.84%。每种尺寸同一进程交错3批、每批100次；批次间存在波动，不是独立启动置信区间，也不代表整模型吞吐提升。未插桩集成A/B和704分辨率仍待验证。

v1补采因误调用pixels.shape()失败，线程栈与CPU复现保留；v2修复后完成，未放宽数值门槛。用户因Codex配额要求本步骤完成后交接，已停止新增实验及CPU自动网页观察进程。完整计划未完成。

证据：[局部输出转换审核](output_conversion_ab_audit_v1.json)；[交接说明](handoff-cn.md)；[交接状态](handoff-state.json)。

## 同一运行的完整kernel账本与关键路径（K0/K1，新增）

新采集：冻结Case D在生产模式（regional torch.compile + reduce-overhead CUDA Graph）下的nsys全进程trace，同一次运行内用NVTX标注request/chunk/step/stage/层（40个block各自是独立CUDA Graph，层级range放在Python调用边界，不改编译代码），VAE worker按叶子模块打range，并在epoch1的chunk0/chunk5记录VAE全部ATen调用元数据。Graph内每个kernel通过cudaGraphLaunch的correlation继承上下文，并用保留的inductor wrapper源码（compile_cache）按launch顺序对齐得到子算子标签。这是诊断trace，不是干净计时；两个带dispatch元数据的VAE chunk另有CPU开销。

### 480x832

覆盖：1,092,983次kernel记录，87.6%的调用和99.9%的稳态GPU时间有NVTX上下文（无上下文部分为初始化/预热/校验/退出）；分组6211个（按角色、上下文、kernel名、launch签名、graph内位置、子算子）。时钟锚点4个，偏移离散122µs。

稳态chunk（epoch1第5–8块）中位数关键路径：DiT请求到返回 677.5 ms（4卡各自GPU忙碌中位 602–603 ms，4卡同时空闲 58.3 ms，其中请求进入到首个kernel约23.7 ms、末尾返回约6.2 ms；每卡未被计算kernel覆盖的NCCL时间中位 44–52 ms）→ 驱动到VAE提交间隙 4.7 ms → IPC提交到worker收到 6.7 ms → H2D 0.53 ms → barrier 2.06 ms → 解码 141.5 ms（4 rank GPU忙碌中位 127–135 ms）→ barrier 0.56 ms → uint8转换 0.16 ms → D2H 1.19 ms → 元数据回传到驱动 15.8 ms；VAE段总墙钟 169.3 ms。Case D各段串行，链路总和即交付间隔；VAE解码约占关键路径17%，DiT约80%。

DiT稳态GPU时间构成（4卡×4间隔合计 9943 ms；按kernel族：gemm 43.6%, attention_core 32.6%, nccl 10.9%, triton_fused 7.5%, gather_index_cat 4.5%, elementwise_reduce 0.8%, attention_prep 0.1%, conv 0.0%）。按子算子（inductor源节点）：

|子算子（源节点）|稳态GPU时间 ms|占比|
|---|---:|---:|
|self-attention：FA3核心|3144.7|31.6%|
|addmm: float_1, modulation, chunk, float_2, layer_norm, scale_msa, add_1, mul|852.0|8.6%|
|addmm: output_parallel_6, input_1, output_parallel_7|721.0|7.3%|
|addmm: output_parallel_6|698.0|7.0%|
|self-attention：paged KV写入/gather（gather_index_cat）|443.7|4.5%|
|_c10d_functional.all_to_all_single: output_parallel, split, value_7, stack, resh|414.8|4.2%|
|addmm: output_parallel_2, input_parallel_1, output_parallel_3|317.4|3.2%|
|_c10d_functional.all_to_all_single: output_5, reshape_3, transpose_3, transpose_|295.2|3.0%|
|addmm: contiguous_12, input_parallel_2, output_parallel_5|293.4|3.0%|
|addmm: float_8, layer_norm_1, float_9, float_10, to_6, output_parallel_4|279.6|2.8%|
|addmm: contiguous_8, input_parallel, output_parallel_1|278.8|2.8%|
|addmm: output_parallel_2|278.1|2.8%|
|addmm: output_parallel_3, camera_features, camera_scale|277.9|2.8%|
|addmm: camera_shift|275.0|2.8%|
|_c10d_functional.all_to_all_single: output_parallel_4, value_float_2, truediv_4,|198.1|2.0%|
|_c10d_functional.all_to_all_single: reshape_10, transpose_8, transpose_9, contig|165.5|1.7%|
|fa3_fwd._flash_attn_forward: _flash_attn_forward_default, key_4, value_9|92.8|0.9%|
|triton_red_fused__to_copy_add_mul_native_layer_norm_split_squeeze_view_9: float_|77.8|0.8%|

VAE稳态GPU时间 2281 ms（按kernel族：conv 44.1%, elementwise_reduce 34.6%, nccl 15.4%, attention_core 4.0%, gather_index_cat 1.7%, gemm 0.2%）；按解码模块前8：vae_frame/* 6.8%；up_blocks.3.resnets.0.conv1/WanDistCausalConv3d 4.9%；up_blocks.3.resnets.0.conv2/WanDistCausalConv3d 4.9%；up_blocks.3.resnets.1.conv1/WanDistCausalConv3d 4.9%；up_blocks.3.resnets.1.conv2/WanDistCausalConv3d 4.9%；up_blocks.3.resnets.2.conv1/WanDistCausalConv3d 4.9%；up_blocks.3.resnets.2.conv2/WanDistCausalConv3d 4.9%；up_blocks.2.upsamplers.0.resample.1/WanDistConv2d 3.7%。

![480x832 账本构成](/ledger_breakdown_480x832.png)

![480x832 逐chunk阶段时间线](/ledger_timeline_480x832.png)

数据：[coverage](/ledger_480x832_coverage_v1.json) · [子算子汇总](/ledger_480x832_summary_v1.json) · [关键路径](/ledger_480x832_critical_path_v1.json) · [graph位置标签](/ledger_480x832_graph_position_labels.json)；完整逐次调用账本 analysis/ledger_480x832_v1/kernel_invocations.jsonl.gz 与 kernel_groups.csv 在本地索引中。

### 704x1280

覆盖：1,092,185次kernel记录，87.5%的调用和100.0%的稳态GPU时间有NVTX上下文（无上下文部分为初始化/预热/校验/退出）；分组6101个（按角色、上下文、kernel名、launch签名、graph内位置、子算子）。时钟锚点4个，偏移离散88µs。

稳态chunk（epoch1第5–8块）中位数关键路径：DiT请求到返回 2138.2 ms（4卡各自GPU忙碌中位 2036–2040 ms，4卡同时空闲 59.1 ms，其中请求进入到首个kernel约29.5 ms、末尾返回约9.9 ms；每卡未被计算kernel覆盖的NCCL时间中位 93–153 ms）→ 驱动到VAE提交间隙 7.4 ms → IPC提交到worker收到 7.1 ms → H2D 0.75 ms → barrier 2.01 ms → 解码 290.0 ms（4 rank GPU忙碌中位 282–282 ms）→ barrier 0.75 ms → uint8转换 0.21 ms → D2H 2.44 ms → 元数据回传到驱动 33.7 ms；VAE段总墙钟 338.0 ms。Case D各段串行，链路总和即交付间隔；VAE解码约占关键路径12%，DiT约86%。

DiT稳态GPU时间构成（4卡×4间隔合计 33525 ms；按kernel族：attention_core 53.0%, gemm 30.1%, nccl 8.7%, triton_fused 4.7%, gather_index_cat 3.0%, elementwise_reduce 0.5%, attention_prep 0.0%, conv 0.0%）。按子算子（inductor源节点）：

|子算子（源节点）|稳态GPU时间 ms|占比|
|---|---:|---:|
|self-attention：FA3核心|17555.8|52.4%|
|addmm: float_1, modulation, chunk, float_2, layer_norm, scale_msa, add_1, mul|1825.2|5.4%|
|addmm: output_parallel_6, input_1, output_parallel_7|1673.0|5.0%|
|addmm: output_parallel_6|1579.8|4.7%|
|self-attention：paged KV写入/gather（gather_index_cat）|1009.1|3.0%|
|_c10d_functional.all_to_all_single: output_parallel, split, value_7, stack, resh|895.0|2.7%|
|_c10d_functional.all_to_all_single: output_5, reshape_3, transpose_3, transpose_|876.2|2.6%|
|addmm: output_parallel_2|861.2|2.6%|
|addmm: output_parallel_2, input_parallel_1, output_parallel_3|804.3|2.4%|
|_c10d_functional.all_to_all_single: output_parallel_4, value_float_2, truediv_4,|724.5|2.2%|
|addmm: contiguous_12, input_parallel_2, output_parallel_5|667.7|2.0%|
|addmm: output_parallel_3, camera_features, camera_scale|665.3|2.0%|
|addmm: contiguous_8, input_parallel, output_parallel_1|653.7|1.9%|
|addmm: float_8, layer_norm_1, float_9, float_10, to_6, output_parallel_4|650.9|1.9%|
|addmm: camera_shift|642.4|1.9%|
|_c10d_functional.all_to_all_single: reshape_10, transpose_8, transpose_9, contig|297.1|0.9%|
|fa3_fwd._flash_attn_forward: _flash_attn_forward_default, key_4, value_9|205.3|0.6%|
|triton_red_fused__to_copy_add_mul_native_layer_norm_split_squeeze_view_9: float_|178.9|0.5%|

VAE稳态GPU时间 4571 ms（按kernel族：conv 48.3%, elementwise_reduce 37.2%, attention_core 9.4%, nccl 3.0%, gather_index_cat 1.9%, gemm 0.2%）；按解码模块前8：vae_frame/* 12.3%；up_blocks.3.resnets.0.conv1/WanDistCausalConv3d 5.4%；up_blocks.3.resnets.0.conv2/WanDistCausalConv3d 5.3%；up_blocks.3.resnets.1.conv2/WanDistCausalConv3d 5.3%；up_blocks.3.resnets.2.conv2/WanDistCausalConv3d 5.3%；up_blocks.3.resnets.1.conv1/WanDistCausalConv3d 5.3%；up_blocks.3.resnets.2.conv1/WanDistCausalConv3d 5.3%；up_blocks.2.upsamplers.0.resample.1/WanDistConv2d 3.9%。

![704x1280 账本构成](/ledger_breakdown_704x1280.png)

![704x1280 逐chunk阶段时间线](/ledger_timeline_704x1280.png)

数据：[coverage](/ledger_704x1280_coverage_v1.json) · [子算子汇总](/ledger_704x1280_summary_v1.json) · [关键路径](/ledger_704x1280_critical_path_v1.json) · [graph位置标签](/ledger_704x1280_graph_position_labels.json)；完整逐次调用账本 analysis/ledger_704x1280_v1/kernel_invocations.jsonl.gz 与 kernel_groups.csv 在本地索引中。

限制：账本时间是nsys记录的kernel时长之和，不是墙钟；graph内位置标签来自wrapper源码对齐，异步NCCL在段内按类别归属；每个kernel的输入shape需结合eager全算子dispatch采集（进行中）按位置对应，尚未逐kernel附上shape；2个VAE chunk带dispatch钩子开销，其解码段偏长（图中chunk 5）。

## Attention有效FLOP利用率曲线（A0，新增）

利用率 = 有效FLOPs /（p50时间 × 规格BF16 dense峰值 989.5 TF/s，来源H200官方页1,979 TFLOPS含稀疏÷2）；另标注同卡经验GEMM上限 788 TF/s（仅一张校准卡）。共1004个点（1004个已审核回放点按配置/角色/阶段/shape/后端聚合，3次重复取中位）。独立算子回放，不是整模型MFU。

|配置|self最高利用率（后端）|cross最高利用率（后端）|
|---|---|---|
|480x832_tp1_sp1|66.1% (native_dense)|52.3% (sdpa_cudnn)|
|480x832_tp1_sp2|67.6% (native_dense)|47.1% (sdpa_cudnn)|
|480x832_tp1_sp4|68.9% (native_dense)|25.9% (sdpa_cudnn)|
|480x832_tp2_sp1|67.2% (native_dense)|47.3% (sdpa_cudnn)|
|480x832_tp2_sp2|67.0% (native_dense)|26.5% (sdpa_cudnn)|
|480x832_tp4_sp1|68.6% (native_dense)|26.8% (vllm_fa2)|
|704x1280_tp1_sp1|69.1% (vllm_fa3)|59.8% (native_dense)|
|704x1280_tp1_sp2|68.4% (native_dense)|54.3% (sdpa_cudnn)|
|704x1280_tp1_sp4|67.0% (native_dense)|47.0% (sdpa_cudnn)|
|704x1280_tp2_sp1|68.7% (native_dense)|54.8% (sdpa_cudnn)|
|704x1280_tp2_sp2|66.1% (native_dense)|46.7% (sdpa_cudnn)|
|704x1280_tp4_sp1|67.2% (native_dense)|47.1% (sdpa_cudnn)|

![self-attention利用率随KV长度](/attention_mfu_overview.png)

各配置分图：[480x832_tp1_sp1](/attention_mfu_480x832_tp1_sp1.png) · [480x832_tp1_sp2](/attention_mfu_480x832_tp1_sp2.png) · [480x832_tp1_sp4](/attention_mfu_480x832_tp1_sp4.png) · [480x832_tp2_sp1](/attention_mfu_480x832_tp2_sp1.png) · [480x832_tp2_sp2](/attention_mfu_480x832_tp2_sp2.png) · [480x832_tp4_sp1](/attention_mfu_480x832_tp4_sp1.png) · [704x1280_tp1_sp1](/attention_mfu_704x1280_tp1_sp1.png) · [704x1280_tp1_sp2](/attention_mfu_704x1280_tp1_sp2.png) · [704x1280_tp1_sp4](/attention_mfu_704x1280_tp1_sp4.png) · [704x1280_tp2_sp1](/attention_mfu_704x1280_tp2_sp1.png) · [704x1280_tp2_sp2](/attention_mfu_704x1280_tp2_sp2.png) · [704x1280_tp4_sp1](/attention_mfu_704x1280_tp4_sp1.png) · [数据点](/attention_mfu_points_v1.json) · [CSV](/attention_mfu_points_v1.csv)

说明：self-attention各配置最高约66–69%（多为native_dense或FA3，含全部有效KV），cross-attention（K=512）利用率低且随并行度下降，是小shape的launch/填充限制而非计算上限；回放GPU未逐卡校准，规格分母为文档值。

## 算子Roofline（G0，DRAM层，新增）

59个已审核真实算子输入（480×832 31个、704×1280 28个）以NCU的dram__bytes读写之和为字节数、标称MAC FLOPs为工作量、NCU累计kernel时长为时间，绘制DRAM层Roofline；族别：GEMM (bf16) 21, conv3d (fp32/tf32) 34, convert/other 4, attention (bf16) 8（convert/copy无标称FLOPs，不进入图，标N/A）。上限线：规格BF16 dense 989.5、TF32 dense 494.5、FP32 67 TF/s 与 DRAM 4.8 TB/s；同卡经验线 788 TF/s / 4.35 TB/s 仅对应一张校准卡。

![算子Roofline](/operator_roofline_v1.png)

[数据点与口径](/operator_roofline_v1.json)

L2层：missing: existing NCU CSVs contain lts__t_sectors_lookup_miss only; lts__t_bytes / l1tex__t_bytes not captured. Re-collection scheduled (operator_l2_traffic_v1).

四类shape曲线（GEMM按M、卷积按工作量、融合elementwise按字节的推导带宽、NCCL按payload；elementwise与NCCL来自同一运行账本，推导带宽不是计数器）：

![shape曲线](/shape_curves_v1.png)

## 逐组瓶颈卡（K2，新增）

对已审核的59个真实算子NCU捕获与8个Attention FA2/FA3捕获，按主导kernel的计数器用公开阈值规则分类（compute_bound: sm or tensor pipe >= 60%; dram_bandwidth_bound: dram >= 60%; l2_bandwidth_bound: lts >= 60% & dram < 60%; partial_wave_tail: sm_active/elapsed < 0.85 & waves <= 1.5; latency_or_occupancy_limited: all < 40% & (warps_active < 25% or waves < 1); else unresolved）。共69张卡：unresolved_no_dominant_limiter 17、compute_bound 29、partial_wave_tail 13、dram_bandwidth_bound 4、l2_bandwidth_bound 6；按阶段：dit → unresolved_no_dominant_limiter 1, compute_bound 14, partial_wave_tail 1, dram_bandwidth_bound 4, l2_bandwidth_bound 1；vae → partial_wave_tail 12, compute_bound 11, l2_bandwidth_bound 5, unresolved_no_dominant_limiter 6；output → unresolved_no_dominant_limiter 4；attention → unresolved_no_dominant_limiter 6, compute_bound 4。每张卡保留原始计数器、触发规则、替代解释与反证实验；unresolved表示没有单一单元≥60%且无占用/尾波特征，需要warp stall分解（已加入排队的L2重采批次）。这是隔离回放的kernel级归因，不是集成流水线归因。

|组|分辨率|主导kernel|分类|sm%|tensor%|dram%|L2%|waves|SM活跃/经过|
|---|---|---|---|---:|---:|---:|---:|---:|---:|
|vae:aten.conv3d.default:{'input': [1, 96, 6, 706, 322], 'weight': [96,|704x1280|sm90_xmma_fprop_implicit_gemm_f32f32_tf3|l2_bandwidth_bound|32|32|10|63|1|1.00|
|vae:aten.conv3d.default:{'input': [1, 192, 6, 354, 162], 'weight': [19|704x1280|sm90_xmma_fprop_implicit_gemm_f32f32_tf3|compute_bound|70|70|9|66|1|0.97|
|vae:aten.conv3d.default:{'input': [1, 96, 6, 482, 210], 'weight': [96,|480x832|sm90_xmma_fprop_implicit_gemm_f32f32_tf3|unresolved_no_dominant_limiter|31|31|10|59|1|0.99|
|vae:aten.conv3d.default:{'input': [1, 96, 6, 706, 322], 'weight': [3, |704x1280|sm80_xmma_fprop_implicit_gemm_indexed_wo|l2_bandwidth_bound|52|20|22|75|11|0.96|
|vae:aten.conv3d.default:{'input': [1, 384, 4, 178, 82], 'weight': [384|704x1280|sm90_xmma_fprop_implicit_gemm_f32f32_tf3|compute_bound|61|61|7|57|1|0.89|
|vae:aten.conv3d.default:{'input': [1, 96, 3, 706, 322], 'weight': [96,|704x1280|sm90_xmma_fprop_implicit_gemm_f32f32_tf3|unresolved_no_dominant_limiter|31|31|10|60|1|0.96|
|vae:aten.conv3d.default:{'input': [1, 192, 6, 242, 106], 'weight': [19|480x832|sm90_xmma_fprop_implicit_gemm_f32f32_tf3|compute_bound|78|78|9|68|1|0.95|
|vae:aten.conv3d.default:{'input': [1, 96, 6, 482, 210], 'weight': [3, |480x832|sm80_xmma_fprop_implicit_gemm_indexed_wo|l2_bandwidth_bound|49|19|21|69|5|0.95|
|dit:aten.linear.default:{'M': 10560, 'N': 5120, 'K': 5120}|704x1280|nvjet_sm90_tst_256x128_64x4_1x2_h_bz_coo|compute_bound|93|93|22|55|1|0.97|
|vae:aten.conv3d.default:{'input': [1, 192, 4, 178, 82], 'weight': [384|704x1280|sm90_xmma_fprop_implicit_gemm_f32f32_tf3|unresolved_no_dominant_limiter|60|60|5|56|1|0.88|
|vae:aten.conv3d.default:{'input': [1, 384, 3, 178, 82], 'weight': [384|704x1280|sm90_xmma_fprop_implicit_gemm_f32f32_tf3|compute_bound|61|61|9|56|1|0.86|
|vae:aten.conv3d.default:{'input': [1, 192, 3, 354, 162], 'weight': [19|704x1280|sm90_xmma_fprop_implicit_gemm_f32f32_tf3|compute_bound|73|73|9|64|1|0.91|
|dit:aten.linear.default:{'M': 2640, 'N': 15360, 'K': 5120}|704x1280|nvjet_sm90_tst_192x208_64x4_2x1_v_bz_coo|compute_bound|93|93|21|48|1|0.96|
|vae:aten.conv3d.default:{'input': [1, 96, 3, 706, 322], 'weight': [3, |704x1280|sm80_xmma_fprop_implicit_gemm_indexed_wo|l2_bandwidth_bound|46|18|18|61|3|0.93|
|vae:aten.conv3d.default:{'input': [1, 384, 4, 122, 54], 'weight': [384|480x832|sm90_xmma_fprop_implicit_gemm_f32f32_tf3|compute_bound|66|66|6|53|1|0.89|
|dit:aten.linear.default:{'M': 2640, 'N': 5120, 'K': 13824}|704x1280|nvjet_sm90_tst_192x192_64x4_1x2_h_bz_coo|compute_bound|92|92|20|44|1|0.93|
|vae:aten.conv3d.default:{'input': [1, 96, 3, 482, 210], 'weight': [96,|480x832|sm90_xmma_fprop_implicit_gemm_f32f32_tf3|l2_bandwidth_bound|30|30|9|60|1|0.95|
|dit:aten.linear.default:{'M': 2640, 'N': 13824, 'K': 5120}|704x1280|nvjet_sm90_tst_128x240_64x4_2x1_v_bz_coo|compute_bound|95|95|21|53|1|0.97|
|vae:aten.conv3d.default:{'input': [1, 192, 3, 178, 82], 'weight': [384|704x1280|sm90_xmma_fprop_implicit_gemm_f32f32_tf3|compute_bound|63|63|8|57|1|0.88|
|vae:aten.conv3d.default:{'input': [1, 384, 4, 176, 80], 'weight': [768|704x1280|sm90_xmma_fprop_implicit_gemm_f32f32_tf3|compute_bound|89|89|33|74|1|0.96|
|dit:aten.linear.default:{'M': 4680, 'N': 5120, 'K': 5120}|480x832|nvjet_sm90_tst_256x144_64x4_2x1_v_bz_coo|compute_bound|94|94|25|53|1|0.98|
|vae:aten.conv3d.default:{'input': [1, 384, 3, 122, 54], 'weight': [384|480x832|sm90_xmma_fprop_implicit_gemm_f32f32_tf3|partial_wave_tail|57|57|7|53|1|0.82|
|vae:aten.conv3d.default:{'input': [1, 96, 3, 482, 210], 'weight': [3, |480x832|sm80_xmma_fprop_implicit_gemm_indexed_wo|partial_wave_tail|33|13|13|42|1|0.79|
|vae:aten.conv3d.default:{'input': [1, 192, 4, 122, 54], 'weight': [384|480x832|sm90_xmma_fprop_implicit_gemm_f32f32_tf3|compute_bound|69|69|5|52|1|0.90|

完整69张卡：[bottleneck_cards_v1.json](/bottleneck_cards_v1.json)（含每个kernel的计数器、规则、替代解释、反证实验和原始NCU目录）。

### 首块 / 缓存增长 / 缓存满 三阶段与请求路径延迟（K0）

同一运行账本按epoch1的chunk分组：首块（chunk 0）、缓存增长（chunk 1–4）、缓存满（chunk 5–9）。“请求进入”= 驱动发出请求到rank0进入_generate_block（引擎/调度/IPC，无GPU工作）；“主机入队完成后GPU尾部”= rank0主机侧完成入队后GPU仍在执行的时间（CUDA Graph使主机先行）；“真实返回延迟”= 最后一个DiT kernel结束到驱动收到latent。

|分辨率|阶段|DiT墙钟 ms|VAE段 ms|请求进入 ms|GPU尾部 ms|真实返回 ms|IPC元数据回传 ms|
|---|---|---:|---:|---:|---:|---:|---:|
|480x832|first_chunk|2377.1|260.7|1888.9|59.3|5.0|10.5|
|480x832|cache_growth|588.8|159.0|28.1|80.5|6.6|13.3|
|480x832|cache_full|678.2|168.6|26.1|98.5|6.0|15.7|
|704x1280|first_chunk|5825.1|337.0|4616.7|201.7|7.7|20.9|
|704x1280|cache_growth|1653.8|336.7|28.0|289.6|8.3|33.8|
|704x1280|cache_full|2141.6|336.9|28.9|399.0|8.7|33.5|

解读：首块的“请求进入”包含会话首块的文本/图像条件编码与缓存初始化（480 约1.9 s、704 约4.6 s，TTFC证据），稳态约26–29 ms；主机在稳态chunk中先于GPU约100 ms（480）/400 ms（704）完成入队，真实返回延迟仅6–8 ms；因此DiT段的关键路径几乎全是GPU执行（含暴露NCCL），主机开销集中在请求进入与元数据回传（15.8/33.7 ms）。

## 账本分组的shape对应与逐组归因覆盖（K1/K2，新增）

把eager全算子dispatch采集（同源码同配置、非同一次运行；40份输出与eager参考逐值一致）的逐调用shape按inductor源节点/顺序附到编译graph的每个位置（融合elementwise kernel的shape按源节点推断并标注inferred），VAE按叶子模块对应；再把每个账本分组对应到瓶颈卡：exact_card=同算子同shape的NCU捕获（含稳态KV长度的Attention），无卡分组用trace推导（shape字节下界/nsys时长）或标communication/unresolved。

|分辨率|稳态GPU时间有shape|计数器归因|通信(trace)|推导归因|未归因|其中“无主导限制单元”的计数器卡占比|
|---|---:|---:|---:|---:|---:|---:|
|480x832|94.7%|68.3%|11.7%|5.7%|14.3%|25.7%|
|704x1280|95.8%|78.3%|6.8%|3.6%|11.4%|46.7%|

“无主导限制单元”主要是稳态KV长度下的FA3 self-attention核心kernel（sm吞吐约55%、DRAM约2%、L2约20%、活跃warp约10%），按现有阈值规则不能单独归为计算或访存受限；已排队的NCU重采会补充warp stall分解、tensor pipe与L2字节，用于最终判定。未归因部分主要是compiled融合kernel（无法用现有算子回放采计数器）、paged KV gather/index kernel和VAE的fmha/elementwise。

480x832 未归因/推导分组前6（稳态GPU ms）：vae vae_frame/* fmha_cutlassF_f32_aligned_32x128_gme 69；dit vllm_omni.ar_diffusion_paged_write_a void at::native::vectorized_gather_k 25；dit vllm_omni.ar_diffusion_paged_write_a void at::native::vectorized_gather_k 25；dit vllm_omni.ar_diffusion_paged_write_a void at::native::vectorized_gather_k 25；dit vllm_omni.ar_diffusion_paged_write_a void at::native::vectorized_gather_k 25；dit vllm_omni.ar_diffusion_paged_write_a void at::native::vectorized_gather_k 24

704x1280 未归因/推导分组前6（稳态GPU ms）：vae vae_frame/* fmha_cutlassF_f32_aligned_32x128_gme 323；vae vae_frame/* fmha_cutlassF_f32_aligned_32x128_gme 107；dit vllm_omni.ar_diffusion_paged_write_a void at::native::vectorized_gather_k 56；dit vllm_omni.ar_diffusion_paged_write_a void at::native::vectorized_gather_k 56；dit vllm_omni.ar_diffusion_paged_write_a void at::native::vectorized_gather_k 56；dit vllm_omni.ar_diffusion_paged_write_a void at::native::vectorized_gather_k 56

数据：[480x832 k2覆盖](/ledger_480x832_k2_coverage_v1.json) · [480x832 逐组归因](/ledger_480x832_bottlenecks_ledger_v1.json) · [704x1280 k2覆盖](/ledger_704x1280_k2_coverage_v1.json) · [704x1280 逐组归因](/ledger_704x1280_bottlenecks_ledger_v1.json)；带shape的分组表 analysis/ledger_<tag>_v1/kernel_groups_with_shapes.csv 在本地索引中。

## 优化迁移矩阵 v2：H1–H5 每个候选的 experiment_id 与状态（H0，新增）

状态只来自磁盘上的结果文件：planned=已定义未跑；running=队列中/运行中；failed_gate=eager数值门槛未过；timed_pending_audit=门槛通过、生产模式计时已跑待审核；bounded_by_ledger=按同一运行账本上界不实施；incompatible_bounded_attempt=有边界移植失败；local_negative=已有局部负结果。计时为6轮生产模式运行的稳态交付间隔（每轮第5–8间隔，共20个），不是1000间隔长尾协议。

|experiment_id|H|候选/单变量|门槛|状态|480 门槛|704 门槛|说明|
|---|---|---|---|---|---|---|---|
|H1a_kv_gather_off|H1|kv_gather_off：LINGBOT_KV_GATHER=0: direct paged FA (block_table) vs gather+FA varlen|threshold|planned|-|-||
|H1b_reduce_relayout|H1|(analysis)：remove pack/unpack clone+transpose kernels around the Ulysses exchange|n/a|bounded_by_ledger|-|-|Same-run ledger: pack/unpack Triton kernels are 480x832 1.67%, 704x1280 0.93% of DiT steady GPU time; upper bo|
|H1c_symm_a2a_transport|H1|(bounded port attempt)：symmetric-memory permute-free all-to-all (a2a_permute.py) for LingBot |exact (data movement only)|incompatible_bounded_attempt|-|-|JIT extension cannot be built in the frozen environment without changing the CUDA toolkit/header set, which wo|
|H2a_camera_injector_chunk_cache|H2|cam_cache：camera injector GEMMs evaluated once per chunk per block (5x fewer), s|exact|planned|-|-|Ledger: camera GEMMs are 11.6% (480) / 8.9% (704) of DiT steady GPU time; 4/5 of that is removable.|
|H2b_norm_rope_modulation_fusion|H2|(analysis)：additional fusion of norm/RoPE/modulation/residual|n/a|bounded_by_ledger|-|-|torch.compile already fuses these into 480x832 7.5%, 704x1280 4.7% of DiT steady GPU time (all Triton fused ke|
|H2c_cudagraph_input_copies|H2|(analysis)：avoid per-block CUDA-graph input staging copies (hidden/camera/expande|n/a|bounded_by_ledger|-|-|Ledger DtoD memcpy in DiT ranks: 480x832 ~80 GB/chunk over 4 ranks, ~11.9 ms copy time per rank per chunk; 704|
|H3a_vae_decoder_compile|H3|vae_compile：torch.compile(vae.decoder) in the four VAE ranks|threshold|planned|-|-||
|H3b_channels_last_local|H3|(prior local A/B)：decoder conv layout|exact|local_negative|-|-|local channels-last A/B including all conversions: exact but 22.42% slower|
|H3c_temporal_conv_layout|H3|(prior local A/B)：decoder conv layout|exact|local_negative|-|-|temporal conv full-cost layout A/B: exact but 14.82% slower|
|H4a_output_conversion_fusion|H4|fused_output：rank0 float32->uint8 conversion: fused Triton kernel|exact|running|未过|未过|local A/B: 20 inputs + 775 rounding edges exact; D2H-inclusive local time -23.76%/-31.84%|
|H4b_pinned_d2h_staging|H4|pinned_d2h：reused pinned staging buffer + non_blocking D2H + pageable clone|exact|planned|-|-||
|H4c_delivery_ipc_encode|H4|(measurement)：n/a: GPU-ready / CPU-ready / cross-process received / encoded timestam|byte-exact frames|planned|-|-|ledger: metadata IPC result latency 15.8 ms (480) / 33.7 ms (704) per chunk; pixel IPC not in frozen path|
|H5a_decode_overlap_case_e|H5|case_e：frozen driver mode E: decode of chunk i overlaps DiT of chunk i+1; sch|exact|planned|-|-||
|H5b_backpressure_cancel|H5|(harness)：bounded queue, slow consumer, cancel, error exit, last-chunk flush|no frame loss/duplication; state cleanup|planned|-|-||

[optimization_transfer_matrix_v2.json](/optimization_transfer_matrix_v2.json) · [H1c 有边界移植尝试记录](/h1_symm_a2a_transport_attempt_v1.json)

## 编译模式跨启动不可重复的分步诊断（NUM，新增）

背景：eager跨启动逐位一致，但生产模式（regional torch.compile + reduce-overhead）两次启动的40份保存输出0/40通过冻结门槛。按约定做有边界的单变量诊断，每步2次独立启动×2 epoch，同一比较脚本与门槛。

|阶段|控制变量|重复性（40份输出）|逐位一致|
|---|---|---|---|
|A_default_mode|LINGBOT_COMPILE_MODE=default（inductor编译，无CUDA Graph树）|0/40 通过|0/40|
|B_reduce_overhead_deterministic|reduce-overhead + CUBLAS_WORKSPACE_CONFIG=:4096:8 + cudnn.benchmark=False/deterministic=True + torch.use_deterministic_algorithms(True, warn_only)|40/40 通过|40/40|

结论（到目前）：CUDA Graph本身不是不可重复的来源（去掉Graph后仍0/40）；加上库级确定性控制后生产模式两次启动逐位一致（40/40 exact）。因此跨启动差异来自cuBLAS工作区/非确定性算法选择等库级非确定性，而非编译或Graph。单一控制拆分（B1 cuBLAS工作区、B2 cuDNN确定性、B3 use_deterministic_algorithms）与该控制的生产模式性能代价计时已排队；若代价可接受，可把确定性设置作为候选验收的生产模式参考，使门槛判定不再依赖eager。

数据：[A 比较](/compile_bisect_A_default_mode_repeat_comparison_v1.json) · [B 比较](/compile_bisect_B_reduce_overhead_deterministic_repeat_comparison_v1.json) · [流程](/compile_bisect_workflow_v1.json)

## 全模型标称FLOPs与模型级MFU（K2/G0，新增）

模型FLOPs来自eager全算子dispatch的逐调用shape（稳态chunk：4次DMD前向+1次commit，4个rank全部求和；self-attention按稳态有效KV长度、非因果、4·Sq·Skv·H·D；线性/卷积按2·MAC，含显式padding），时间来自同一运行账本的关键路径中位数（诊断trace，非干净计时）。分母：DiT用规格BF16 dense峰值989.5 TF/s，VAE卷积为TF32 implicit GEMM（xmma f32f32_tf32f32 kernel）用TF32 dense峰值494.5 TF/s。这是“模型MFU”口径；硬件执行FLOPs/Tensor Core活跃度/逐kernel利用率是不同指标。

|分辨率|DiT标称TFLOP/chunk（linear/self/cross）|DiT墙钟MFU|DiT按GPU忙碌MFU|VAE卷积TFLOP/chunk|VAE解码MFU(TF32)|交付间隔|端到端MFU(BF16参考)|每卡达到TF/s|
|---|---|---:|---:|---:|---:|---:|---:|---:|
|480x832|1315（767/538/9.8）|49.0%|55.2%|40.5|14.5%|852 ms (14.09 FPS)|40.2%|398|
|704x1280|4494（1731/2741/22.1）|53.1%|55.7%|91.5|15.9%|2484 ms (4.83 FPS)|46.6%|462|

解读：DiT在稳态chunk内按墙钟约49–53%的规格BF16 dense峰值，其中self-attention核心kernel在NCU隔离回放中sm吞吐约55%（见瓶颈卡）、GEMM多为计算受限；4卡同时空闲与暴露NCCL约占DiT墙钟10–11%。VAE解码为FP32/TF32卷积，模型级MFU约14–16%，与逐kernel归因（小grid尾波、layout转换、elementwise）一致。704×1280的DiT FLOPs为480的3.4倍（attention 5.1倍），交付间隔2.9倍，端到端MFU反而更高，说明480×832更受launch/通信/非计算段影响。低MFU本身不是瓶颈结论，各段归因见上文。

数据：[480x832](/model_mfu_480x832_v1.json) · [704x1280](/model_mfu_704x1280_v1.json)

## 候选整模型A/B结果（G1，持续更新）

门槛：480 eager 2 epoch的40份保存输出对eager参考（exact或冻结容差，按候选契约）；计时：生产模式（reduce-overhead）6轮=1预热+5测量轮、每轮第5–8个CPU-ready交付间隔，原版/候选/原版/候选交错在同一分配上运行；比值为中位数比与percentile bootstrap 95% CI。不是1000间隔长尾协议，且这里的原版两次启动之比给出噪声下限。

|候选|分辨率|eager门槛|原版稳态间隔中位 ms|候选稳态间隔中位 ms|候选/原版 中位比 [95% CI]|判定|
|---|---|---|---:|---:|---|---|
|H4a 输出转换融合|480x832|40/40 过门槛，40/40 逐位|826.0 (n=40)|835.1 (n=40)|1.0110 [1.0051, 1.0162]|no resolvable difference (effect within CI or within run-to-run spread of either arm)；原版两次启动比 1.0033|
|H4a 输出转换融合|704x1280|40/40 过门槛，40/40 逐位|-|-|-|-|