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)。
已知负结果(保留,不重试同口径):
- 两种编译模式(default、reduce-overhead)跨启动重复均 0/40 通过门槛;eager 跨启动逐位一致,是候选判定的唯一参考。
- cuDNN cross 与 SDPA Flash cross 候选:局部通过,整模型 0/40 通过门槛;不接受。
- 两个含全部转换成本的 VAE channels-last 局部 A/B:逐值一致但分别慢 22.42%、14.82%;不接受。
- 长尾 6 次运行仅 epoch0/1 保存完整输出,其余不可补回;最终验收长尾将重跑并保存全部输出。
- 704×1280 稳态约 4.89 FPS,与 12 FPS 目标差约 2.46 倍;480×832 约 14.4 FPS 已高于 12。
唯一局部正结果:输出转换 Triton 融合,20 份真实输入与 775 个取整边界值逐字节一致,含 D2H 局部耗时降低 23.76%/31.84%;尚未集成整模型验收。
当前可用结论
首轮固定报告包已在2026-09-10 08:40(北京时间)交付,以下为持续补充的在线版本;完整研究计划仍未全部完成。
- Attention:12个配置的代表上下文共尝试1008个回放点,1004个通过审核,4个SDPA math点显存不足;并非所有模型层完整覆盖。
- 硬件证据:59种真实线性/卷积/转换代表输入共183次NCU kernel记录(含一项调度异常备注),以及同卡FA2/FA3对照已审核;原始输入、计数器和输出证据可复用。
- 同卡参考:计算与连续读写带宽校准已完成,12个算子纳入经验参考图;这不是整模型MFU或理论硬件峰值。
- 优化验收:cuDNN cross候选未通过整模型误差门槛;含全部转换成本的VAE局部channels-last候选反而慢约22.42%,不认领加速。SDPA Flash cross整模型候选40份输出也全部未通过门槛;新参考40份逐值一致,失败证据与单次耗时见末节。
- 性能边界:现有交付计时主要到CPU-ready uint8;编码、网络、客户端播放、背压与取消仍未建立完整证据。
- 数据限制:六次长尾运行保存了逐次计时/哈希,但后续轮次未保留完整输出;已缺失的完整输出不能从哈希恢复。所有新实验继续保存独立目录和校验清单。
后文保留历史失败、诊断与各次补充,均按测量范围解释;旧阶段的限制如已解决,以对应更新章节为准。
首轮固定版本报告包:下载tar.gz(10.54MB) · 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。
- latent相对L2范围:0.029937–0.403371。
- 归一化像素RMSE范围:0.017613–0.291942。
这是原始模型跨启动控制,数值差异原因尚未定位;不能据此接受优化或放宽正确性门槛。逐文件误差与哈希保存在analysis/tail_704_r1_r2_saved_output_comparison.json。
704×1280两次原始模型运行的保存输出对比
已比较40个保存张量,40个未通过冻结数值门槛。范围仅为epoch 0、1的latent与像素输出,不能推广到全部250个测量session。
- latent相对L2范围:0.036485–0.309353。
- 归一化像素RMSE范围:0.019445–0.212082。
这是原始模型跨启动控制,数值差异原因尚未定位;不能据此接受优化或放宽正确性门槛。逐文件误差与哈希保存在analysis/tail_704_r0_r1_saved_output_comparison.json。
第二轮长尾样本的时间戳分解
第二轮1000个稳态间隔也通过非负与分段求和检查。最大样本位于epoch=32、chunk=8,交付间隔1010.481 ms。其中latent-ready到发出解码为180.323 ms,该段全样本中位数为4.817 ms;该样本DiT区间675.661 ms,decoder计算区间130.184 ms。
本轮最大样本的额外间隔与首轮所在位置不同,不能将所有长尾归因于同一个阶段。这里仅定位主机标记之间的间隔,没有CPU采样和逐事件关联,不能确认具体函数、调度抢占或通信原因。
下载第二轮1000个间隔分解CSV。证据: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 矢量图 SVG


下载 480x832_tp1_sp2 矢量图 SVG


下载 480x832_tp1_sp4 矢量图 SVG


下载 480x832_tp2_sp1 矢量图 SVG


下载 480x832_tp2_sp2 矢量图 SVG


下载 480x832_tp4_sp1 矢量图 SVG


下载 704x1280_tp1_sp1 矢量图 SVG


下载 704x1280_tp1_sp2 矢量图 SVG


下载 704x1280_tp1_sp4 矢量图 SVG


下载 704x1280_tp2_sp1 矢量图 SVG


下载 704x1280_tp2_sp2 矢量图 SVG


下载 704x1280_tp4_sp1 矢量图 SVG


下载 mfu_480x832_tp1_sp1 矢量图 SVG

下载 mfu_480x832_tp1_sp2 矢量图 SVG

下载 mfu_480x832_tp1_sp4 矢量图 SVG

下载 mfu_480x832_tp2_sp1 矢量图 SVG

下载 mfu_480x832_tp2_sp2 矢量图 SVG

下载 mfu_480x832_tp4_sp1 矢量图 SVG

下载 mfu_704x1280_tp1_sp1 矢量图 SVG

下载 mfu_704x1280_tp1_sp2 矢量图 SVG

下载 mfu_704x1280_tp1_sp4 矢量图 SVG

下载 mfu_704x1280_tp2_sp1 矢量图 SVG

下载 mfu_704x1280_tp2_sp2 矢量图 SVG

下载 mfu_704x1280_tp4_sp1 矢量图 SVG

下载 mfu_overview 矢量图 SVG

下载 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重复的全部已保存输出逐位一致。

| 比较(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固定报告包,在线报告追加更新,不修改原包。
只读本地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和采集回放影响仍需保留。

逐算子数据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误差。整模型数值失败结论保持。

候选轨迹独立审核 · 候选输出不变检查 · 完整配对数据 · 矢量图
小工作集带宽:尚不能作为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时间占比,右侧是累计时间的对数坐标;都不表示模型墙钟占比或可恢复加速。

较大卷积的计算占比约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%。


数据: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%。


数据: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) |

各配置分图: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 仅对应一张校准卡。

数据点与口径
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