LingBot 全流水线分析:首轮报告(持续补充)

生成时间:2026-09-10T16:50:03.582468+08:00

当前结论只覆盖已核验的数据。完整计划尚未完成;所有局部 Attention 结果均不代表整模型优化已经通过。

首页摘要:16项验收状态与证据入口

验收表生成时间 2026-09-10T08:45:13.764688+00:00(UTC);full_plan_complete=False。状态只来自 analysis/plan_acceptance_audit.json,本摘要不会把部分完成升级为通过。完成路线见 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 · 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_launch_resources_summary.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_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 · fa3_interface_dispatch_comparison.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 · local_gemm_calibration_audit.json · operator_work_traffic.json
G1 整模型优化验收 未通过 cuDNN cross候选40份输出均未通过门槛;尚无接受的整模型加速。 新SDPA Flash候选的6个实际实现输入已完成独立数值和kernel审核;该结果不能替代整模型验收。 本次检查未确认模型任务仍存活,不能仅凭历史状态认定在运行。 新候选保存输出数值检查未通过;单次A/B… cross_candidate_eager_r0_comparison.json · execution_control_pilot_timings.json · cross_sdpa_flash_impl_audit.json
H0 参考优化适用性 部分完成 已做本地源码映射;所列外部PR/RFC状态及本地提交祖先关系已追加核验,完整编译融合及transport运行关联仍缺。 optimization_transfer_matrix.json · h3_external_reference_review.json · h3_reference_local_ancestry.json
H1 Attention与通信布局消融 未完成 独立后端回放不能替代通信布局的单变量A/B。 optimization_transfer_matrix.json
H2 DiT融合与局部数据构造消融 未完成 尚无通过数值门槛的融合增量收益验证。 optimization_transfer_matrix.json
H3 解码算子与空间并行消融 未完成 新增含全部转换成本的局部布局A/B:数值逐值一致但慢约22.42%;仅单输入局部反证,完整解码优化消融仍缺。 新增时间维卷积全成本布局A/B同样逐值一致但慢14.82%;布局转换51.9%时间占比未直接转化为收益。 optimization_transfer_matrix.json · missing_vae_components_v1.json · vae_layout_temporal_ab_audit_v1.json
H4 输出与传输优化 未完成 已有uint8路径与交接计时,尚无受控输出优化及编码后媒体验收。 已补充六阶段生产者/消费者表:冻结路径只回传元数据,不含像素IPC;CPU-ready不能替代编码或客户端交付。输出融合真实输入补采已提交,尚未验收。 新增固定FFV1无损CPU编码往返,3份视频和351帧逐字节审… clean_handoff_summary.json · optimization_transfer_matrix.json · output_boundary_inventory_v1.json
H5 流水线重叠、背压与取消 未完成 时间线重叠不等于优化收益;慢消费者、取消和组合调度验收尚缺。 worker_chunk_activity.json · optimization_transfer_matrix.json
TAIL 重复运行与尾延迟 部分完成 两分辨率各3次、每次1000间隔;新增连续会话块重采样显示部分P99区间敏感,数值控制与未来稳定性仍未完成。 raw/tail_timing_audit_workflow_v1.json · clean_handoff_summary.json · tail_six_launch_hash_repeatability.json
NUM 数值可重复性与候选门槛 部分完成 eager两次40份输出逐位一致;两种编译模式重复及cuDNN候选均失败,原因未最终定位。 新参考与历史40份逐值一致;原版和候选两条轨迹各160次同输入比较及FP64抽样通过,诊断最终输出各40份与对应未插桩运行逐值一致。跨轨迹Q从第2次出现差异,仍缺完整逐层因果证明。 新增分… execution_control_eager_repeat_comparison.json · execution_error_progression.json · cross_sdpa_fresh_reference_repeat_v2.json
RAW 全量原始数据与复现 部分完成 已存在文件有索引、源快照与复算脚本;长尾后续轮次未保存完整输出,无法补回。 delivery_archive_audit_20260910T0020Z.json · REPRODUCE_CN.md · raw_data_index.json
WEB 首轮报告与证据访问 首版已交付/持续更新 中文报告及看板已发布,82个报告内资源链接通过内容核对;不代表完整计划完成。 REPORT_CN.md · public_report_asset_audit.json

已知负结果(保留,不重试同口径):

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

当前可用结论

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

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

首轮固定版本报告包下载tar.gz(10.54MB) · 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来源说明已保留。
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%/推导14%/未归因6%;704x1280 计数器归因78%/通信7%/推导10%/未归因5%。稳态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。此表是审核时点快照;引用文件之后可能继续更新,快照中的哈希用于识别当时版本。首版报告交付不等于所有要求通过。

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

此前整模型诊断记录中,最后一个具名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名称。证据: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 · 第三轮分布图。证据: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:P99的进程内session分组bootstrap 95%区间为836.430–842.137 ms;不是独立启动之间的置信区间。

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

全部稳态间隔及累计分布

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

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

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

全部稳态间隔及累计分布

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

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

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

全部稳态间隔及累计分布

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

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

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

全部稳态间隔及累计分布

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

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

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

全部稳态间隔及累计分布

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

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

下载tail_original_704x1280_r2_v1全部逐块计时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。

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

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

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

这是原始模型跨启动控制,数值差异原因尚未定位;不能据此接受优化或放宽正确性门槛。逐文件误差与哈希保存在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。证据: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。证据: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。证据: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。证据: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。证据: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。原始证据及检查: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

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。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局部胜出不改变整模型数值门槛未通过的事实,不能直接推荐替换。

局部最快接口与缺失配置

图中星号表示与至少一个竞争者的重复范围重叠;灰格包括缺失或因OOM未满足七接口比较条件的上下文;OOM明确标注。每个单元格固定本配置的真实输入,不比较不同TP/SP行的端到端收益。下载逐点结果 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。复核证据: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
rebased 262 101350 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完整稳态名称排名

拷贝方向 调用次数 累计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完整稳态名称排名

拷贝方向 调用次数 累计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结果,不证明整模型数值等价。

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

480x832_tp1_sp1 Attention 实测散点图

下载 480x832_tp1_sp1 矢量图 SVG

480x832_tp1_sp1 同输入耗时比值热力图

480x832_tp1_sp2 Attention 实测散点图

下载 480x832_tp1_sp2 矢量图 SVG

480x832_tp1_sp2 同输入耗时比值热力图

480x832_tp1_sp4 Attention 实测散点图

下载 480x832_tp1_sp4 矢量图 SVG

480x832_tp1_sp4 同输入耗时比值热力图

480x832_tp2_sp1 Attention 实测散点图

下载 480x832_tp2_sp1 矢量图 SVG

480x832_tp2_sp1 同输入耗时比值热力图

480x832_tp2_sp2 Attention 实测散点图

下载 480x832_tp2_sp2 矢量图 SVG

480x832_tp2_sp2 同输入耗时比值热力图

480x832_tp4_sp1 Attention 实测散点图

下载 480x832_tp4_sp1 矢量图 SVG

480x832_tp4_sp1 同输入耗时比值热力图

704x1280_tp1_sp1 Attention 实测散点图

下载 704x1280_tp1_sp1 矢量图 SVG

704x1280_tp1_sp1 同输入耗时比值热力图

704x1280_tp1_sp2 Attention 实测散点图

下载 704x1280_tp1_sp2 矢量图 SVG

704x1280_tp1_sp2 同输入耗时比值热力图

704x1280_tp1_sp4 Attention 实测散点图

下载 704x1280_tp1_sp4 矢量图 SVG

704x1280_tp1_sp4 同输入耗时比值热力图

704x1280_tp2_sp1 Attention 实测散点图

下载 704x1280_tp2_sp1 矢量图 SVG

704x1280_tp2_sp1 同输入耗时比值热力图

704x1280_tp2_sp2 Attention 实测散点图

下载 704x1280_tp2_sp2 矢量图 SVG

704x1280_tp2_sp2 同输入耗时比值热力图

704x1280_tp4_sp1 Attention 实测散点图

下载 704x1280_tp4_sp1 矢量图 SVG

704x1280_tp4_sp1 同输入耗时比值热力图

mfu_480x832_tp1_sp1 Attention 实测散点图

下载 mfu_480x832_tp1_sp1 矢量图 SVG

mfu_480x832_tp1_sp2 Attention 实测散点图

下载 mfu_480x832_tp1_sp2 矢量图 SVG

mfu_480x832_tp1_sp4 Attention 实测散点图

下载 mfu_480x832_tp1_sp4 矢量图 SVG

mfu_480x832_tp2_sp1 Attention 实测散点图

下载 mfu_480x832_tp2_sp1 矢量图 SVG

mfu_480x832_tp2_sp2 Attention 实测散点图

下载 mfu_480x832_tp2_sp2 矢量图 SVG

mfu_480x832_tp4_sp1 Attention 实测散点图

下载 mfu_480x832_tp4_sp1 矢量图 SVG

mfu_704x1280_tp1_sp1 Attention 实测散点图

下载 mfu_704x1280_tp1_sp1 矢量图 SVG

mfu_704x1280_tp1_sp2 Attention 实测散点图

下载 mfu_704x1280_tp1_sp2 矢量图 SVG

mfu_704x1280_tp1_sp4 Attention 实测散点图

下载 mfu_704x1280_tp1_sp4 矢量图 SVG

mfu_704x1280_tp2_sp1 Attention 实测散点图

下载 mfu_704x1280_tp2_sp1 矢量图 SVG

mfu_704x1280_tp2_sp2 Attention 实测散点图

下载 mfu_704x1280_tp2_sp2 矢量图 SVG

mfu_704x1280_tp4_sp1 Attention 实测散点图

下载 mfu_704x1280_tp4_sp1 矢量图 SVG

mfu_overview Attention 实测散点图

下载 mfu_overview 矢量图 SVG

winners Attention 实测散点图

下载 winners 矢量图 SVG

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

下载全部逐组图表数据 CSV

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

下载比值及参考后端 JSON

归档索引扫描时间:2026-09-10T06:13:26.519937+00:00;记录33055个文件、约399.31GB。这是扫描时点快照,活动文件仍需在任务结束后重新封存。索引位于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分组;时钟与末块存在性通过不等于逐层调用完整,也不意味着瓶颈已确定。旧记录的限制独立保留。

下载调用覆盖清单 · 下载逐组待归因清单。生成脚本: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。证据: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输出误差

比较(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 · 下载矢量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。完整数据: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。原始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 · 下载分组统计及最慢样本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。复算脚本:scripts/audit_operator_signature_ambiguity.py。

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

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

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

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

H3官方文章的第6.3节仍将原始benchmark bundle标为待发布。外部文章数字继续仅作背景,不用来补本机缺失计数器、拟合MFU或承诺端到端收益。

下载外部响应元数据及SHA256 · 下载本地提交祖先核对。原始网页/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。

权限验证审核 · 三项真实算子审核与指标

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项完整计数器审核与指标

同卡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对照原始指标、数值误差与哈希

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

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

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

独立算子工作量与VAE局部布局对照

逐算子数据JSON · CSV · 矢量图

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次。是一次进程内的相关样本,不是三次独立启动,也不是整模型尾延迟验收。反例不证明所有布局优化均无效;跨层保持布局或权重预打包需重新定义公平边界并另测。

详细硬件计数器 · 详细采集审核与调度备注 · 局部A/B审核

本机连续读写带宽校准

已完成同一物理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,不把两块卡数据直接当作同卡硬上界。

三次启动审核与原始文件哈希 · NCU流量核对 · 物理GPU身份审核

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

三次新的独立启动已完成: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

同卡代表算子的实测参考图

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官方规格区分SXM/NVL,列出的Tensor Core数字带稀疏性脚注。不能把该稀疏数值直接当成本实验稠密运算的MFU分母,也不凭L20X驱动名称选择峰值。

同卡计算审核 · 计算CSV · 参考图完整数据与排除原因 · 矢量图 · 原始PCI映射审核

新的整模型候选: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。排队或运行状态不是验证成功。

六个实际实现验证的审核

新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。

完整输出误差 · 来源与计时审核 · 日志名称修正依据

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

源码快照索引及逐块误差

后续真实输入诊断: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次同输入独立审核 · 诊断输出一致性

候选自身轨迹诊断: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差异

候选轨迹独立审核 · 候选输出不变检查 · 完整配对数据 · 矢量图

小工作集带宽:尚不能作为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计时审核 · 原生NCU指标和单位

小工作集提交方式对照

同一进程、同一输入和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及校验

已记录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 · 代表采集原始行复算

逐调用有效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份输出已保存。

长度与覆盖审核 · 40份输出一致性 · 逐调用账本(gzip)

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

复用六次既有运行,每次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

长尾区间对连续会话块长的敏感性

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

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

全部4种块长的数据与哈希 · 独立核对记录 · 矢量图

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

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映射 · 算子记录审核 · Attention长度审核 · 最终输出一致性

代表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

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

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

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

逐签名工作量、排序与复核

新增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项硬件指标与原始报告校验 · 形状/工作量/DRAM流量 · 真实快照审核 · 模型输出一致性 · 补齐后结构对应

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

按明确的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分类及时间复核

时间维卷积布局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%,不接受为优化收益。

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

完整原始计时/输出审核

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数据、调用数与源文件校验

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或放宽数值门槛。

分辨率工作量对照 · 704逐chunk工作量 · 704结构覆盖 · 来源/尺寸/输出审核 · 全部模型输出比较 · 同次工作量与物理映射

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计时不可乘调用数当模型墙钟时间。

计数器/输出审核与调度备注 · 工作量/流量和逐项备注 · 补齐后的704结构覆盖 · 真实输入审核 · 模型输出一致性

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

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

704×1280 VAE各真实形状的计算与转换构成

较大卷积的计算占比约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分组与源证据 · 矢量图

短 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 输出并非逐位一致。局部容差通过不代表整模型输出通过。

证据:汇总数据;原始目录: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。证据:输入匹配及计数器数据

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

冻结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在该冻结基线未执行,服务边界仍需独立实验。

证据:输出边界与逻辑尺寸清单。输出融合补采和局部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、逐帧哈希、命令和日志均已保留。

证据:媒体往返审核;原始目录: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自动网页观察进程。完整计划未完成。

证据:局部输出转换审核交接说明交接状态

同一运行的完整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 账本构成

480x832 逐chunk阶段时间线

数据:coverage · 子算子汇总 · 关键路径 · graph位置标签;完整逐次调用账本 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 账本构成

704x1280 逐chunk阶段时间线

数据:coverage · 子算子汇总 · 关键路径 · graph位置标签;完整逐次调用账本 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长度

各配置分图:480x832_tp1_sp1 · 480x832_tp1_sp2 · 480x832_tp1_sp4 · 480x832_tp2_sp1 · 480x832_tp2_sp2 · 480x832_tp4_sp1 · 704x1280_tp1_sp1 · 704x1280_tp1_sp2 · 704x1280_tp1_sp4 · 704x1280_tp2_sp1 · 704x1280_tp2_sp2 · 704x1280_tp4_sp1 · 数据点 · 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

数据点与口径

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).

逐组瓶颈卡(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(含每个kernel的计数器、规则、替代解释、反证实验和原始NCU目录)。

账本分组的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% 13.8% 6.2% 25.7%
704x1280 95.8% 78.3% 6.8% 9.6% 5.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覆盖 · 480x832 逐组归因 · 704x1280 k2覆盖 · 704x1280 逐组归因;带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
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 · H1c 有边界移植尝试记录

编译模式跨启动不可重复的分步诊断(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 比较 · B 比较 · 流程

全模型标称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 · 704x1280