Record TP2 prefill serving-path verdict

This commit is contained in:
2026-07-23 18:08:32 +08:00
parent cf610003ed
commit 4f22688bfd
9 changed files with 2013 additions and 5 deletions

View File

@@ -1,3 +1,4 @@
fleet-artifacts/ fleet-artifacts/
fleet-state/ fleet-state/
replay/
remote-outputs/ remote-outputs/

View File

@@ -0,0 +1,113 @@
#!/usr/bin/env python3
"""Compare the measured-prefill-MoE TP2 replays with structured baseline."""
from __future__ import annotations
import importlib.util
import json
import sys
from pathlib import Path
ROOT = Path(__file__).resolve().parent
REPO = ROOT.parents[1]
ATTN = REPO / "runs/frontier-attn-structured-v0"
S3_REAL = REPO / "runs/frontier-s3-real-v0"
CELLS = {
"tp2_rho0p0025": "frontier-s3-real-full-r0p0025-tp2-t*",
"tp2_rho0p005": "frontier-s3-real-full-r0p005-tp2-t*",
}
def load_analysis_module():
path = ATTN / "analyze_trace_verdict.py"
spec = importlib.util.spec_from_file_location("attention_verdict", path)
module = importlib.util.module_from_spec(spec)
sys.path.insert(0, str(ATTN))
spec.loader.exec_module(module)
return module
def main() -> None:
analysis = load_analysis_module()
cells = {}
for label, pattern in CELLS.items():
real_trials = analysis.load_real_trials(pattern)
structured = analysis.load_sim(ATTN / "replay" / label)
moe_corrected = analysis.load_sim(ROOT / "replay" / label)
structured_biases = [
analysis.distribution_bias(trial, structured)
for trial in real_trials
]
corrected_biases = [
analysis.distribution_bias(trial, moe_corrected)
for trial in real_trials
]
cells[label] = {
"structured_attention": {
"trialwise_distribution_bias": structured_biases,
"trialwise_distribution_bias_summary": (
analysis.aggregate_trial_bias(structured_biases)
),
"legacy_pooled_distribution_bias": (
analysis.legacy_pooled_bias(real_trials, structured)
),
"waiting_p99_ms": analysis.waiting_p99(structured),
},
"structured_attention_plus_prefill_moe": {
"trialwise_distribution_bias": corrected_biases,
"trialwise_distribution_bias_summary": (
analysis.aggregate_trial_bias(corrected_biases)
),
"legacy_pooled_distribution_bias": (
analysis.legacy_pooled_bias(real_trials, moe_corrected)
),
"paired_relative_error": [
analysis.paired_relative_error(trial, moe_corrected)
for trial in real_trials
],
"waiting_p99_ms": analysis.waiting_p99(moe_corrected),
},
}
low = cells["tp2_rho0p0025"]
before = low["structured_attention"]["legacy_pooled_distribution_bias"]
after = low["structured_attention_plus_prefill_moe"][
"legacy_pooled_distribution_bias"
]
checked = [
(metric, quantile)
for metric in ("ttft", "e2e")
for quantile in ("mean", "p50", "p99")
]
gates = {
"subcritical_waiting_below_1s": (
low["structured_attention_plus_prefill_moe"]["waiting_p99_ms"]
< 1000
),
"subcritical_ttft_e2e_abs_bias_not_worse": all(
abs(after[metric][quantile])
<= abs(before[metric][quantile]) + 0.01
for metric, quantile in checked
),
"subcritical_mean_ttft_abs_bias_improves_3pp": (
abs(before["ttft"]["mean"]) - abs(after["ttft"]["mean"])
>= 0.03
),
}
payload = {
"schema": "frontier-tp2-prefill-serving-replay-verdict.v1",
"cells": cells,
"gates": gates,
"decision": (
"keep_tp2_prefill_moe_correction"
if all(gates.values())
else "reject_global_scale_and_fit_shape_conditioned_curve"
),
}
output = ROOT / "results/replay-verdict.json"
output.write_text(json.dumps(payload, indent=2, sort_keys=True) + "\n")
print(json.dumps({"gates": gates, "decision": payload["decision"]}, indent=2))
if __name__ == "__main__":
main()

View File

@@ -0,0 +1,80 @@
#!/usr/bin/env python3
"""Combine the frozen simulator entry audit with the serving trace smoke."""
from __future__ import annotations
import json
from pathlib import Path
ROOT = Path(__file__).resolve().parent
ENTRY = ROOT / "results/entry-audit.json"
SMOKE = ROOT / "results/serving-smoke.json"
def main() -> None:
entry = json.loads(ENTRY.read_text())
smoke = json.loads(SMOKE.read_text())
simulator = entry["cells"]["tp2"]
critical = smoke["critical_rank"]
sim_total = float(simulator["total_ms"])
real_total = float(simulator["real_total_ms"])
sim_moe = float(simulator["moe_grouped_gemm_ms"])
serving_total = float(critical["execute_wall_ms"])
serving_moe = float(critical["components_ms"]["moe"])
total_residual = real_total - sim_total
moe_residual = serving_moe - sim_moe
counterfactual_total = sim_total + moe_residual
payload = {
"schema": "frontier-tp2-prefill-serving-smoke-verdict.v1",
"contract": {
"simulator_sample": "mean of first nine q8192/ctx0 single-request chunks",
"serving_sample": "longest execute window on critical TP rank",
"moe_mapping": (
"standalone moe_grouped_gemm and serving MoE both include "
"expert prepare/finalize plus expert GEMMs"
),
"real_anchor_ms": real_total,
},
"measurements_ms": {
"simulator_total": sim_total,
"serving_execute_wall": serving_total,
"real_anchor": real_total,
"simulator_moe_grouped_gemm": sim_moe,
"serving_moe": serving_moe,
"simulator_non_moe": sim_total - sim_moe,
"serving_non_moe": serving_total - serving_moe,
},
"counterfactual": {
"total_residual_ms": total_residual,
"moe_residual_ms": moe_residual,
"moe_residual_fraction": moe_residual / total_residual,
"moe_scale": serving_moe / sim_moe,
"moe_only_counterfactual_total_ms": counterfactual_total,
"moe_only_counterfactual_bias": (
counterfactual_total - real_total
)
/ real_total,
},
"gates": {
"serving_reproduces_anchor_within_5pct": (
abs(serving_total - real_total) / real_total <= 0.05
),
"moe_explains_at_least_70pct": moe_residual / total_residual >= 0.70,
"moe_only_counterfactual_within_5pct": (
abs(counterfactual_total - real_total) / real_total <= 0.05
),
},
}
payload["decision"] = (
"inject_tp2_prefill_moe_and_replay"
if all(payload["gates"].values())
else "stop_moe_injection_and_profile_whole_layer"
)
output = ROOT / "results/serving-smoke-verdict.json"
output.write_text(json.dumps(payload, indent=2, sort_keys=True) + "\n")
print(json.dumps(payload, indent=2, sort_keys=True))
if __name__ == "__main__":
main()

View File

@@ -1,6 +1,6 @@
# 实验 EXP-TP2-PREFILL-SERVINGTP2 base-prefill residual 是否来自 serving-path MoE # 实验 EXP-TP2-PREFILL-SERVINGTP2 base-prefill residual 是否来自 serving-path MoE
> **状态:** GPU smoke harness 已冻结,待远端执行 > **状态:** 完成;机制 PASSglobal constant injection FAIL
> >
> Parent campaign[`../frontier-simulator-gap-campaign-v0/README.md`](../frontier-simulator-gap-campaign-v0/README.md) > Parent campaign[`../frontier-simulator-gap-campaign-v0/README.md`](../frontier-simulator-gap-campaign-v0/README.md)
@@ -84,10 +84,26 @@
real=`410 ms`bias=`13.11%`TP4=`231.33 vs 231 ms`。TP2 real=`410 ms`bias=`13.11%`TP4=`231.33 vs 231 ms`。TP2
component ledger MoE=`171.12 ms`,若单独解释 residual 需增至约 component ledger MoE=`171.12 ms`,若单独解释 residual 需增至约
`224.9 ms`+31.4%)。 `224.9 ms`+31.4%)。
- **观察事实** GPU - **观察事实**
- TP2 q8192 profile execute=`408.19 ms`,冻结 real anchor=`410 ms`
两个 TP rank 分别为 `408.19/407.34 ms`
- critical-rank attention=`99.67 ms`sim attention execution=`100.24 ms`
serving MoE=`214.52 ms`sim MoE=`171.12 ms`
- MoE delta=`43.41 ms`,解释 total residual `80.75%`MoE-only
counterfactual total=`399.65 ms`(相对 real `2.52%`)。三个事前 smoke
gate 全部 PASS
- 但把 q8k ratio `1.25366×` 用作所有 TP2 prefill shape 的常数 scale
ρ=0.0025 TTFT mean `4.54%→+5.23%`p99
`7.67%→+0.89%`E2E mean `+12.70%→+15.55%`ρ=0.005
TTFT mean `7.08%→+3.13%`E2E mean `+5.02%→+9.14%`
- **Fleet preflight2026-07-23** dash1--dash4 均为 8×H2032 张卡 - **Fleet preflight2026-07-23** dash1--dash4 均为 8×H2032 张卡
`memory.used=0 MiB``utilization=0%` compute processdry-run 选择 `memory.used=0 MiB``utilization=0%` compute processdry-run 选择
`dash1:[0,1]`正式 job pin dash1 `dash1:[0,1]`正式 job pin dash1
- **含义** GPU - **含义** TP2 base residual 的主要机制确实是 serving-path MoE而不是
- **Claim update** unchanged attention real-anchor 噪声但单个 q8k 点不能外推成全 prefill-domain
- **下一步** probe fleet运行 TP2 serving-path prefill smoke constant calibration
- **Claim update** TP2 prefill MoE 尚有明显工程可优化 gap得到支持
用一个 TP2 常数 scale 即可修复 trace fidelity被否定
- **下一步** 将后续工程项收窄为 TP2 token/routing-conditioned serving
MoE curve至少 q2k/q4k/q8k 与真实 routing allocation不合入当前
global scalecampaign 继续实验 3

View File

@@ -0,0 +1,525 @@
{
"cc_cache": "/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/cc-cache",
"cells": {
"tp1_mns16": {
"argv": [
"/usr/bin/python3",
"/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/run_frontier_with_curves.py",
"--simulation_mode",
"online",
"--sys_arch",
"co-location",
"--cc_backend_config_type",
"vidur",
"--cluster_config_num_replicas",
"1",
"--cluster_scheduler_config_type",
"sticky_round_robin",
"--replica_config_model_name",
"qwen3-a3b-30b-moe",
"--replica_config_device",
"h20",
"--replica_config_network_device",
"h20_dgx",
"--replica_config_attn_tensor_parallel_size",
"1",
"--replica_config_attn_data_parallel_size",
"1",
"--replica_config_moe_tensor_parallel_size",
"1",
"--replica_config_moe_expert_parallel_size",
"1",
"--replica_config_num_pipeline_stages",
"1",
"--replica_scheduler_config_type",
"vllm_v1",
"--decode_cuda_graph_mode",
"piecewise",
"--vllm_v1_scheduler_config_batch_size_cap",
"16",
"--vllm_v1_scheduler_config_max_tokens_in_batch",
"8192",
"--vllm_v1_scheduler_config_long_prefill_token_threshold",
"0",
"--vllm_v1_scheduler_config_block_size",
"16",
"--vllm_v1_scheduler_config_num_blocks_mode",
"explicit",
"--vllm_v1_scheduler_config_gpu_memory_utilization",
"0.92",
"--vllm_v1_scheduler_config_non_kv_cache_overhead_bytes",
"0",
"--request_generator_config_type",
"trace_replay",
"--trace_request_generator_config_trace_file",
"/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/joint-r2/inputs/tp1-frontier.csv",
"--trace_request_generator_config_max_tokens",
"40960",
"--metrics_config_output_dir",
"/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/joint-r2/sim/tp1_mns16/metrics",
"--metrics_config_run_id",
"joint_tp1_mns16",
"--metrics_config_write_metrics",
"--metrics_config_store_request_metrics",
"--metrics_config_store_batch_metrics",
"--metrics_config_store_token_completion_metrics",
"--metrics_config_store_utilization_metrics",
"--no-metrics_config_store_plots",
"--no-metrics_config_enable_chrome_trace",
"--no-metrics_config_write_json_trace",
"--metrics_config_store_frontier_stage_batch_ledger",
"--no-random_forrest_execution_time_predictor_config_enable_dummy_mode",
"--random_forrest_execution_time_predictor_config_linear_op_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/profile-v4-trace-final/linear_op.csv",
"--random_forrest_execution_time_predictor_config_atten_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/profile-v4-trace-final/attention.csv",
"--random_forrest_execution_time_predictor_config_moe_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/profile-v4-trace-final/moe.csv",
"--random_forrest_execution_time_predictor_config_linear_op_kernel_only_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/frozen-kernel-only/linear_op.csv",
"--random_forrest_execution_time_predictor_config_atten_kernel_only_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/frozen-kernel-only/attention.csv",
"--random_forrest_execution_time_predictor_config_moe_kernel_only_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/frozen-kernel-only/moe.csv",
"--random_forrest_execution_time_predictor_config_prediction_max_prefill_chunk_size",
"8192",
"--random_forrest_execution_time_predictor_config_prediction_max_batch_size",
"32",
"--random_forrest_execution_time_predictor_config_prediction_max_tokens_per_request",
"40960",
"--random_forrest_execution_time_predictor_config_no_cache",
"--random_forrest_execution_time_predictor_config_skip_cpu_overhead_modeling",
"--vllm_v1_scheduler_config_num_blocks",
"20128",
"--vllm_v1_scheduler_config_enable_chunked_prefill",
"--random_forrest_execution_time_predictor_config_num_training_job_threads",
"4",
"--cudagraph_capture_sizes",
"1",
"2",
"4",
"8",
"16",
"24",
"32",
"--vidur_cc_backend_config_all_reduce_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/measured-allreduce.csv",
"--vidur_cc_backend_config_cache_dir",
"/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/cc-cache",
"--vidur_cc_backend_config_k_fold_cv_splits",
"6",
"--vidur_cc_backend_config_num_training_job_threads",
"1",
"--metrics_config_cache_dir",
"/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/model-cache"
],
"log": "/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/joint-r2/logs/tp1_mns16.log",
"source_command": "/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-lo-fixed-pd-cells/sim/fixed-pd/runs/tp1_mns16/tp1/command.json",
"source_command_sha256": "a9815797b1601bf6f6cdf0269e84acb376a84945609e338868dc8347aab650e6",
"usage": "/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/joint-r2/usage/tp1_mns16.json"
},
"tp2_mns16": {
"argv": [
"/usr/bin/python3",
"/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/run_frontier_with_curves.py",
"--simulation_mode",
"online",
"--sys_arch",
"co-location",
"--cc_backend_config_type",
"vidur",
"--cluster_config_num_replicas",
"1",
"--cluster_scheduler_config_type",
"sticky_round_robin",
"--replica_config_model_name",
"qwen3-a3b-30b-moe",
"--replica_config_device",
"h20",
"--replica_config_network_device",
"h20_dgx",
"--replica_config_attn_tensor_parallel_size",
"2",
"--replica_config_attn_data_parallel_size",
"1",
"--replica_config_moe_tensor_parallel_size",
"2",
"--replica_config_moe_expert_parallel_size",
"1",
"--replica_config_num_pipeline_stages",
"1",
"--replica_scheduler_config_type",
"vllm_v1",
"--decode_cuda_graph_mode",
"piecewise",
"--vllm_v1_scheduler_config_batch_size_cap",
"16",
"--vllm_v1_scheduler_config_max_tokens_in_batch",
"8192",
"--vllm_v1_scheduler_config_long_prefill_token_threshold",
"0",
"--vllm_v1_scheduler_config_block_size",
"16",
"--vllm_v1_scheduler_config_num_blocks_mode",
"explicit",
"--vllm_v1_scheduler_config_gpu_memory_utilization",
"0.92",
"--vllm_v1_scheduler_config_non_kv_cache_overhead_bytes",
"0",
"--request_generator_config_type",
"trace_replay",
"--trace_request_generator_config_trace_file",
"/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/joint-r2/inputs/tp2-frontier.csv",
"--trace_request_generator_config_max_tokens",
"40960",
"--metrics_config_output_dir",
"/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/joint-r2/sim/tp2_mns16/metrics",
"--metrics_config_run_id",
"joint_tp2_mns16",
"--metrics_config_write_metrics",
"--metrics_config_store_request_metrics",
"--metrics_config_store_batch_metrics",
"--metrics_config_store_token_completion_metrics",
"--metrics_config_store_utilization_metrics",
"--no-metrics_config_store_plots",
"--no-metrics_config_enable_chrome_trace",
"--no-metrics_config_write_json_trace",
"--metrics_config_store_frontier_stage_batch_ledger",
"--no-random_forrest_execution_time_predictor_config_enable_dummy_mode",
"--random_forrest_execution_time_predictor_config_linear_op_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/profile-v4-trace-final/linear_op.csv",
"--random_forrest_execution_time_predictor_config_atten_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/profile-v4-trace-final/attention.csv",
"--random_forrest_execution_time_predictor_config_moe_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/profile-v4-trace-final/moe.csv",
"--random_forrest_execution_time_predictor_config_linear_op_kernel_only_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/frozen-kernel-only/linear_op.csv",
"--random_forrest_execution_time_predictor_config_atten_kernel_only_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/frozen-kernel-only/attention.csv",
"--random_forrest_execution_time_predictor_config_moe_kernel_only_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/frozen-kernel-only/moe.csv",
"--random_forrest_execution_time_predictor_config_prediction_max_prefill_chunk_size",
"8192",
"--random_forrest_execution_time_predictor_config_prediction_max_batch_size",
"32",
"--random_forrest_execution_time_predictor_config_prediction_max_tokens_per_request",
"40960",
"--random_forrest_execution_time_predictor_config_no_cache",
"--random_forrest_execution_time_predictor_config_skip_cpu_overhead_modeling",
"--vllm_v1_scheduler_config_num_blocks",
"76620",
"--vllm_v1_scheduler_config_enable_chunked_prefill",
"--random_forrest_execution_time_predictor_config_num_training_job_threads",
"4",
"--cudagraph_capture_sizes",
"1",
"2",
"4",
"8",
"16",
"24",
"32",
"--vidur_cc_backend_config_all_reduce_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/measured-allreduce.csv",
"--vidur_cc_backend_config_cache_dir",
"/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/cc-cache",
"--vidur_cc_backend_config_k_fold_cv_splits",
"6",
"--vidur_cc_backend_config_num_training_job_threads",
"1",
"--metrics_config_cache_dir",
"/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/model-cache"
],
"log": "/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/joint-r2/logs/tp2_mns16.log",
"source_command": "/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-lo-fixed-pd-cells/sim/fixed-pd/runs/tp2_mns16/tp2/command.json",
"source_command_sha256": "61788a8810be301c9dbc006624aa19b6a932bc44d341b836861087833cffc3df",
"usage": "/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/joint-r2/usage/tp2_mns16.json"
},
"tp4_mns16": {
"argv": [
"/usr/bin/python3",
"/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/run_frontier_with_curves.py",
"--simulation_mode",
"online",
"--sys_arch",
"co-location",
"--cc_backend_config_type",
"vidur",
"--cluster_config_num_replicas",
"1",
"--cluster_scheduler_config_type",
"sticky_round_robin",
"--replica_config_model_name",
"qwen3-a3b-30b-moe",
"--replica_config_device",
"h20",
"--replica_config_network_device",
"h20_dgx",
"--replica_config_attn_tensor_parallel_size",
"4",
"--replica_config_attn_data_parallel_size",
"1",
"--replica_config_moe_tensor_parallel_size",
"4",
"--replica_config_moe_expert_parallel_size",
"1",
"--replica_config_num_pipeline_stages",
"1",
"--replica_scheduler_config_type",
"vllm_v1",
"--decode_cuda_graph_mode",
"piecewise",
"--vllm_v1_scheduler_config_batch_size_cap",
"16",
"--vllm_v1_scheduler_config_max_tokens_in_batch",
"8192",
"--vllm_v1_scheduler_config_long_prefill_token_threshold",
"0",
"--vllm_v1_scheduler_config_block_size",
"16",
"--vllm_v1_scheduler_config_num_blocks_mode",
"explicit",
"--vllm_v1_scheduler_config_gpu_memory_utilization",
"0.92",
"--vllm_v1_scheduler_config_non_kv_cache_overhead_bytes",
"0",
"--request_generator_config_type",
"trace_replay",
"--trace_request_generator_config_trace_file",
"/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/joint-r2/inputs/tp4-frontier.csv",
"--trace_request_generator_config_max_tokens",
"40960",
"--metrics_config_output_dir",
"/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/joint-r2/sim/tp4_mns16/metrics",
"--metrics_config_run_id",
"joint_tp4_mns16",
"--metrics_config_write_metrics",
"--metrics_config_store_request_metrics",
"--metrics_config_store_batch_metrics",
"--metrics_config_store_token_completion_metrics",
"--metrics_config_store_utilization_metrics",
"--no-metrics_config_store_plots",
"--no-metrics_config_enable_chrome_trace",
"--no-metrics_config_write_json_trace",
"--metrics_config_store_frontier_stage_batch_ledger",
"--no-random_forrest_execution_time_predictor_config_enable_dummy_mode",
"--random_forrest_execution_time_predictor_config_linear_op_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/profile-v4-trace-final/linear_op.csv",
"--random_forrest_execution_time_predictor_config_atten_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/profile-v4-trace-final/attention.csv",
"--random_forrest_execution_time_predictor_config_moe_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/profile-v4-trace-final/moe.csv",
"--random_forrest_execution_time_predictor_config_linear_op_kernel_only_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/frozen-kernel-only/linear_op.csv",
"--random_forrest_execution_time_predictor_config_atten_kernel_only_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/frozen-kernel-only/attention.csv",
"--random_forrest_execution_time_predictor_config_moe_kernel_only_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/frozen-kernel-only/moe.csv",
"--random_forrest_execution_time_predictor_config_prediction_max_prefill_chunk_size",
"8192",
"--random_forrest_execution_time_predictor_config_prediction_max_batch_size",
"32",
"--random_forrest_execution_time_predictor_config_prediction_max_tokens_per_request",
"40960",
"--random_forrest_execution_time_predictor_config_no_cache",
"--random_forrest_execution_time_predictor_config_skip_cpu_overhead_modeling",
"--vllm_v1_scheduler_config_num_blocks",
"191882",
"--vllm_v1_scheduler_config_enable_chunked_prefill",
"--random_forrest_execution_time_predictor_config_num_training_job_threads",
"4",
"--cudagraph_capture_sizes",
"1",
"2",
"4",
"8",
"16",
"24",
"32",
"--vidur_cc_backend_config_all_reduce_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/measured-allreduce.csv",
"--vidur_cc_backend_config_cache_dir",
"/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/cc-cache",
"--vidur_cc_backend_config_k_fold_cv_splits",
"6",
"--vidur_cc_backend_config_num_training_job_threads",
"1",
"--metrics_config_cache_dir",
"/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/model-cache"
],
"log": "/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/joint-r2/logs/tp4_mns16.log",
"source_command": "/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-lo-fixed-pd-cells/sim/fixed-pd/runs/tp4_mns16/tp4/command.json",
"source_command_sha256": "9bbcf10446336ba5885193f391dd628cd18ff64d91a51ebcd463ffd24be95532",
"usage": "/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/joint-r2/usage/tp4_mns16.json"
},
"tp4_mns32": {
"argv": [
"/usr/bin/python3",
"/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/run_frontier_with_curves.py",
"--simulation_mode",
"online",
"--sys_arch",
"co-location",
"--cc_backend_config_type",
"vidur",
"--cluster_config_num_replicas",
"1",
"--cluster_scheduler_config_type",
"sticky_round_robin",
"--replica_config_model_name",
"qwen3-a3b-30b-moe",
"--replica_config_device",
"h20",
"--replica_config_network_device",
"h20_dgx",
"--replica_config_attn_tensor_parallel_size",
"4",
"--replica_config_attn_data_parallel_size",
"1",
"--replica_config_moe_tensor_parallel_size",
"4",
"--replica_config_moe_expert_parallel_size",
"1",
"--replica_config_num_pipeline_stages",
"1",
"--replica_scheduler_config_type",
"vllm_v1",
"--decode_cuda_graph_mode",
"piecewise",
"--vllm_v1_scheduler_config_batch_size_cap",
"32",
"--vllm_v1_scheduler_config_max_tokens_in_batch",
"8192",
"--vllm_v1_scheduler_config_long_prefill_token_threshold",
"0",
"--vllm_v1_scheduler_config_block_size",
"16",
"--vllm_v1_scheduler_config_num_blocks_mode",
"explicit",
"--vllm_v1_scheduler_config_gpu_memory_utilization",
"0.92",
"--vllm_v1_scheduler_config_non_kv_cache_overhead_bytes",
"0",
"--request_generator_config_type",
"trace_replay",
"--trace_request_generator_config_trace_file",
"/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/joint-r2/inputs/tp4-frontier.csv",
"--trace_request_generator_config_max_tokens",
"40960",
"--metrics_config_output_dir",
"/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/joint-r2/sim/tp4_mns32/metrics",
"--metrics_config_run_id",
"joint_tp4_mns32",
"--metrics_config_write_metrics",
"--metrics_config_store_request_metrics",
"--metrics_config_store_batch_metrics",
"--metrics_config_store_token_completion_metrics",
"--metrics_config_store_utilization_metrics",
"--no-metrics_config_store_plots",
"--no-metrics_config_enable_chrome_trace",
"--no-metrics_config_write_json_trace",
"--metrics_config_store_frontier_stage_batch_ledger",
"--no-random_forrest_execution_time_predictor_config_enable_dummy_mode",
"--random_forrest_execution_time_predictor_config_linear_op_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/profile-v4-trace-final/linear_op.csv",
"--random_forrest_execution_time_predictor_config_atten_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/profile-v4-trace-final/attention.csv",
"--random_forrest_execution_time_predictor_config_moe_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/profile-v4-trace-final/moe.csv",
"--random_forrest_execution_time_predictor_config_linear_op_kernel_only_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/frozen-kernel-only/linear_op.csv",
"--random_forrest_execution_time_predictor_config_atten_kernel_only_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/frozen-kernel-only/attention.csv",
"--random_forrest_execution_time_predictor_config_moe_kernel_only_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/frozen-kernel-only/moe.csv",
"--random_forrest_execution_time_predictor_config_prediction_max_prefill_chunk_size",
"8192",
"--random_forrest_execution_time_predictor_config_prediction_max_batch_size",
"64",
"--random_forrest_execution_time_predictor_config_prediction_max_tokens_per_request",
"40960",
"--random_forrest_execution_time_predictor_config_no_cache",
"--random_forrest_execution_time_predictor_config_skip_cpu_overhead_modeling",
"--vllm_v1_scheduler_config_num_blocks",
"191786",
"--vllm_v1_scheduler_config_enable_chunked_prefill",
"--random_forrest_execution_time_predictor_config_num_training_job_threads",
"4",
"--cudagraph_capture_sizes",
"1",
"2",
"4",
"8",
"16",
"24",
"32",
"40",
"48",
"56",
"64",
"--vidur_cc_backend_config_all_reduce_input_file",
"/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-profiles/measured-allreduce.csv",
"--vidur_cc_backend_config_cache_dir",
"/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/cc-cache",
"--vidur_cc_backend_config_k_fold_cv_splits",
"6",
"--vidur_cc_backend_config_num_training_job_threads",
"1",
"--metrics_config_cache_dir",
"/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/model-cache"
],
"log": "/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/joint-r2/logs/tp4_mns32.log",
"source_command": "/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-lo-fixed-pd-cells/sim/fixed-pd/runs/tp4_mns32/tp4/command.json",
"source_command_sha256": "fbc7dee55590b415ed1cde8072de835ed155a0c20ba0eb305c3cb22aa8065a51",
"usage": "/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/joint-r2/usage/tp4_mns32.json"
}
},
"collective_curve": "/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/results/collective-curve.json",
"collective_curve_sha256": "f9543649d4ea78f08240bf1284ab74083aa5cf5671ed47e386047f1453300b36",
"collective_curve_variant": "drop_mean",
"frontier_checkout": "/tmp/frontier-attn-structured-v0",
"frontier_commit": "1f8900a4ac64e45754b03d0aa7c1dddab65785cf",
"mode": "joint",
"model_cache": "/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/model-cache",
"moe_curve": "/home/gahow/phd/aituner/runs/frontier-fused-moe-profile-v0/results/fused-moe-curve.json",
"moe_curve_sha256": "b94d65d9d581adefcc6c14ed4920cce6a1136f74f1014737dc3e4249bc8250d2",
"python": "/usr/bin/python3",
"python_dependency_roots": [
"/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/python-deps",
"/home/gahow/.cache/uv/archive-v0/-_kzErLcPO5nASZFX8b9k",
"/home/gahow/.cache/uv/archive-v0/FbaBs_QJ9QKEbQ9V_4aIR",
"/home/gahow/.cache/uv/archive-v0/fuHsGXD0Lv_UjFC8yI4-7",
"/home/gahow/.cache/uv/archive-v0/jFGdqQLpB1eopfm9VxT3j",
"/home/gahow/.cache/uv/archive-v0/YWW6ExSJuPVvv4-qYQTin",
"/home/gahow/.cache/uv/archive-v0/3_qxZ5Ll-EpVAGZfbksfe"
],
"traces": {
"1": {
"first_arrival_s": 0.0,
"last_arrival_s": 595.348837209302,
"requests": 129,
"source_request_metrics": "/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-lo-fixed-pd-cells/sim/fixed-pd/runs/tp1_mns16/tp1/metrics/qwen3_a3b_30b_moe/online_serving/qwen30_trace_tp1_mns16_tp1/request_metrics.csv",
"source_sha256": "0b82e09644a5884fcd10d894b68495daefdabb32b770146c2f9ece37b8469f4f",
"trace": "/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/joint-r2/inputs/tp1-frontier.csv",
"trace_sha256": "59dd8996ff879ef94330004104dfdf515b791bce4036576eccc93290e9206dad"
},
"2": {
"first_arrival_s": 0.0,
"last_arrival_s": 297.674418604651,
"requests": 129,
"source_request_metrics": "/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-lo-fixed-pd-cells/sim/fixed-pd/runs/tp2_mns16/tp2/metrics/qwen3_a3b_30b_moe/online_serving/qwen30_trace_tp2_mns16_tp2/request_metrics.csv",
"source_sha256": "33983081bb20dd5e2053e9e3d13def8732e958150c9b47a8609ba345123f2316",
"trace": "/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/joint-r2/inputs/tp2-frontier.csv",
"trace_sha256": "64fc077b38274a76a8279884ac4115836cd1157c95119c64fabac50d81124f69"
},
"4": {
"first_arrival_s": 0.0,
"last_arrival_s": 148.837209302326,
"requests": 129,
"source_request_metrics": "/home/gahow/phd/aituner/runs/frontier-split-rootcause-v0/frozen-inputs/q30-lo-fixed-pd-cells/sim/fixed-pd/runs/tp4_mns16/tp4/metrics/qwen3_a3b_30b_moe/online_serving/qwen30_trace_tp4_mns16_tp4/request_metrics.csv",
"source_sha256": "b36cd383c07b546d2c1f2fac754d5dbb92efd6880316b4233a7aef9fa1115a36",
"trace": "/home/gahow/phd/aituner/runs/frontier-collective-joint-v0/counterfactual/joint-r2/inputs/tp4-frontier.csv",
"trace_sha256": "adb3d6f3932a44c86c3d9e7cf1e57739594e8c48d19a26b5aa54b37dce0e0c19"
}
}
}

View File

@@ -0,0 +1,631 @@
{
"cells": {
"tp2_rho0p0025": {
"structured_attention": {
"legacy_pooled_distribution_bias": {
"e2e": {
"mean": 0.12696146154030977,
"p50": 0.14796251184179712,
"p90": 0.12767235986604622,
"p99": -0.03965040123988098
},
"tpot": {
"mean": 0.15073709978689778,
"p50": 0.18527285925405948,
"p90": 0.07863996768588354,
"p99": 0.1661995397125918
},
"ttft": {
"mean": -0.045433491316312524,
"p50": -0.0875990583320057,
"p90": -0.1142200855839675,
"p99": -0.07674087497700505
}
},
"trialwise_distribution_bias": [
{
"e2e": {
"mean": 0.12073678040141796,
"p50": 0.14420262806746403,
"p90": 0.11915070375608593,
"p99": -0.022788557992646938
},
"tpot": {
"mean": 0.14386158526258666,
"p50": 0.1787588067010671,
"p90": 0.071419448886029,
"p99": 0.17844287277146695
},
"ttft": {
"mean": -0.044844475267623835,
"p50": -0.08906525404741798,
"p90": -0.11412907883561804,
"p99": -0.07227786537885415
}
},
{
"e2e": {
"mean": 0.13325567383154438,
"p50": 0.15920445684060733,
"p90": 0.12953600816854335,
"p99": -0.009483068681846096
},
"tpot": {
"mean": 0.15769576871598226,
"p50": 0.196444372264794,
"p90": 0.08604019319402298,
"p99": 0.19311481061458743
},
"ttft": {
"mean": -0.046021781355297144,
"p50": -0.08612813517467069,
"p90": -0.11510233828040589,
"p99": -0.07592149320892366
}
}
],
"trialwise_distribution_bias_summary": {
"e2e": {
"mean": {
"max": 0.13325567383154438,
"mean": 0.12699622711648118,
"min": 0.12073678040141796
},
"p50": {
"max": 0.15920445684060733,
"mean": 0.15170354245403567,
"min": 0.14420262806746403
},
"p90": {
"max": 0.12953600816854335,
"mean": 0.12434335596231463,
"min": 0.11915070375608593
},
"p99": {
"max": -0.009483068681846096,
"mean": -0.016135813337246518,
"min": -0.022788557992646938
}
},
"tpot": {
"mean": {
"max": 0.15769576871598226,
"mean": 0.15077867698928446,
"min": 0.14386158526258666
},
"p50": {
"max": 0.196444372264794,
"mean": 0.18760158948293054,
"min": 0.1787588067010671
},
"p90": {
"max": 0.08604019319402298,
"mean": 0.07872982104002599,
"min": 0.071419448886029
},
"p99": {
"max": 0.19311481061458743,
"mean": 0.1857788416930272,
"min": 0.17844287277146695
}
},
"ttft": {
"mean": {
"max": -0.044844475267623835,
"mean": -0.045433128311460486,
"min": -0.046021781355297144
},
"p50": {
"max": -0.08612813517467069,
"mean": -0.08759669461104433,
"min": -0.08906525404741798
},
"p90": {
"max": -0.11412907883561804,
"mean": -0.11461570855801197,
"min": -0.11510233828040589
},
"p99": {
"max": -0.07227786537885415,
"mean": -0.07409967929388891,
"min": -0.07592149320892366
}
}
},
"waiting_p99_ms": 777.5911879793948
},
"structured_attention_plus_prefill_moe": {
"legacy_pooled_distribution_bias": {
"e2e": {
"mean": 0.15546889966298288,
"p50": 0.160521791081281,
"p90": 0.14949973986858572,
"p99": 0.063501869119007
},
"tpot": {
"mean": 0.1725158165141388,
"p50": 0.19940581724371484,
"p90": 0.1197438720681557,
"p99": 0.16591937069699617
},
"ttft": {
"mean": 0.052253838015785536,
"p50": 0.011133339751522397,
"p90": -0.05603209782959642,
"p99": 0.008949154409837737
}
},
"paired_relative_error": [
{
"e2e": {
"mean": 0.19703071945786896,
"p50": 0.16651188156434055,
"p90": 0.3164763187383559,
"p99": 0.828609074267838
},
"tpot": {
"mean": 0.19036780497451272,
"p50": 0.17682334345090506,
"p90": 0.25886898169288464,
"p99": 0.7758893841183759
},
"ttft": {
"mean": 0.45065657187067015,
"p50": 0.10957981115881457,
"p90": 1.0421699185581517,
"p99": 5.507647228917042
}
},
{
"e2e": {
"mean": 0.20969149313321356,
"p50": 0.17863746376866904,
"p90": 0.32770775080455516,
"p99": 0.8245897427826396
},
"tpot": {
"mean": 0.20429023428176118,
"p50": 0.19259101506782267,
"p90": 0.2812533323043662,
"p99": 0.7385621812554619
},
"ttft": {
"mean": 0.4457497323795849,
"p50": 0.10597514413253763,
"p90": 1.0162200916250081,
"p99": 5.496662131059387
}
}
],
"trialwise_distribution_bias": [
{
"e2e": {
"mean": 0.14908675997874046,
"p50": 0.1567207723136509,
"p90": 0.14081313742067922,
"p99": 0.08217486261364416
},
"tpot": {
"mean": 0.16551017680032373,
"p50": 0.1928140924228096,
"p90": 0.11224819981281003,
"p99": 0.1781597624045383
},
"ttft": {
"mean": 0.052903132111322415,
"p50": 0.00950848459985634,
"p90": -0.05593511273439991,
"p99": 0.013826387288597744
}
},
{
"e2e": {
"mean": 0.16192232934855333,
"p50": 0.1718867285340943,
"p90": 0.15139946120181186,
"p99": 0.09690951004812934
},
"tpot": {
"mean": 0.17960618440333612,
"p50": 0.21071053715514645,
"p90": 0.1274260991438306,
"p99": 0.19282817544573386
},
"ttft": {
"mean": 0.051605344227498705,
"p50": 0.012763433909333126,
"p90": -0.05697230680643937,
"p99": 0.009844585085406704
}
}
],
"trialwise_distribution_bias_summary": {
"e2e": {
"mean": {
"max": 0.16192232934855333,
"mean": 0.15550454466364688,
"min": 0.14908675997874046
},
"p50": {
"max": 0.1718867285340943,
"mean": 0.1643037504238726,
"min": 0.1567207723136509
},
"p90": {
"max": 0.15139946120181186,
"mean": 0.14610629931124552,
"min": 0.14081313742067922
},
"p99": {
"max": 0.09690951004812934,
"mean": 0.08954218633088676,
"min": 0.08217486261364416
}
},
"tpot": {
"mean": {
"max": 0.17960618440333612,
"mean": 0.17255818060182992,
"min": 0.16551017680032373
},
"p50": {
"max": 0.21071053715514645,
"mean": 0.20176231478897805,
"min": 0.1928140924228096
},
"p90": {
"max": 0.1274260991438306,
"mean": 0.11983714947832032,
"min": 0.11224819981281003
},
"p99": {
"max": 0.19282817544573386,
"mean": 0.1854939689251361,
"min": 0.1781597624045383
}
},
"ttft": {
"mean": {
"max": 0.052903132111322415,
"mean": 0.05225423816941056,
"min": 0.051605344227498705
},
"p50": {
"max": 0.012763433909333126,
"mean": 0.011135959254594732,
"min": 0.00950848459985634
},
"p90": {
"max": -0.05593511273439991,
"mean": -0.056453709770419636,
"min": -0.05697230680643937
},
"p99": {
"max": 0.013826387288597744,
"mean": 0.011835486187002224,
"min": 0.009844585085406704
}
}
},
"waiting_p99_ms": 872.1814900223759
}
},
"tp2_rho0p005": {
"structured_attention": {
"legacy_pooled_distribution_bias": {
"e2e": {
"mean": 0.050164236759908075,
"p50": 0.0912493994395978,
"p90": -0.0004979191794830456,
"p99": -0.0014369075488634014
},
"tpot": {
"mean": 0.05457336599027876,
"p50": 0.0872512146277895,
"p90": 0.01132352325010614,
"p99": -0.05948329733667248
},
"ttft": {
"mean": -0.07082500585057615,
"p50": -0.1284088888361358,
"p90": -0.111136573344055,
"p99": -0.09012556161973535
}
},
"trialwise_distribution_bias": [
{
"e2e": {
"mean": 0.04462683161260099,
"p50": 0.0878889042422101,
"p90": -0.006843953727060337,
"p99": -0.0036641600308731643
},
"tpot": {
"mean": 0.048585629713612626,
"p50": 0.08234652581183928,
"p90": 0.003101343693636994,
"p99": -0.05952673172238401
},
"ttft": {
"mean": -0.07113138559444453,
"p50": -0.1294796632798084,
"p90": -0.11063584772248519,
"p99": -0.08982438463776078
}
},
{
"e2e": {
"mean": 0.05576066061144686,
"p50": 0.09176507061640418,
"p90": 0.003187406757202862,
"p99": 0.004813545234813467
},
"tpot": {
"mean": 0.060629878515092876,
"p50": 0.09375745746657616,
"p90": 0.019190079238111446,
"p99": -0.05840233314016777
},
"ttft": {
"mean": -0.07051842392629604,
"p50": -0.12771160295843229,
"p90": -0.11011498510890981,
"p99": -0.08861845103384182
}
}
],
"trialwise_distribution_bias_summary": {
"e2e": {
"mean": {
"max": 0.05576066061144686,
"mean": 0.05019374611202393,
"min": 0.04462683161260099
},
"p50": {
"max": 0.09176507061640418,
"mean": 0.08982698742930714,
"min": 0.0878889042422101
},
"p90": {
"max": 0.003187406757202862,
"mean": -0.0018282734849287376,
"min": -0.006843953727060337
},
"p99": {
"max": 0.004813545234813467,
"mean": 0.0005746926019701514,
"min": -0.0036641600308731643
}
},
"tpot": {
"mean": {
"max": 0.060629878515092876,
"mean": 0.05460775411435275,
"min": 0.048585629713612626
},
"p50": {
"max": 0.09375745746657616,
"mean": 0.08805199163920771,
"min": 0.08234652581183928
},
"p90": {
"max": 0.019190079238111446,
"mean": 0.01114571146587422,
"min": 0.003101343693636994
},
"p99": {
"max": -0.05840233314016777,
"mean": -0.05896453243127589,
"min": -0.05952673172238401
}
},
"ttft": {
"mean": {
"max": -0.07051842392629604,
"mean": -0.07082490476037029,
"min": -0.07113138559444453
},
"p50": {
"max": -0.12771160295843229,
"mean": -0.12859563311912034,
"min": -0.1294796632798084
},
"p90": {
"max": -0.11011498510890981,
"mean": -0.11037541641569751,
"min": -0.11063584772248519
},
"p99": {
"max": -0.08861845103384182,
"mean": -0.0892214178358013,
"min": -0.08982438463776078
}
}
},
"waiting_p99_ms": 1166.1196883368564
},
"structured_attention_plus_prefill_moe": {
"legacy_pooled_distribution_bias": {
"e2e": {
"mean": 0.09135599701015996,
"p50": 0.12135156828958114,
"p90": 0.046904152670232756,
"p99": 0.05267198541612694
},
"tpot": {
"mean": 0.09686033389444083,
"p50": 0.10920022438478016,
"p90": 0.055959658792935615,
"p99": 0.03647201411457465
},
"ttft": {
"mean": 0.03133826519177879,
"p50": -0.002404965655955949,
"p90": -0.03671038522787783,
"p99": -0.003047423826180572
}
},
"paired_relative_error": [
{
"e2e": {
"mean": 0.1373684769346999,
"p50": 0.10813242136980938,
"p90": 0.27973507165610256,
"p99": 0.6892985502092539
},
"tpot": {
"mean": 0.12750409894388354,
"p50": 0.10486952080742089,
"p90": 0.2347358304319833,
"p99": 0.7326960496792728
},
"ttft": {
"mean": 0.48959967245129044,
"p50": 0.11864693915356177,
"p90": 1.3203954670852478,
"p99": 5.486884563706639
}
},
{
"e2e": {
"mean": 0.15015620590606898,
"p50": 0.11828746639583118,
"p90": 0.29344565738935435,
"p99": 0.7209062198737742
},
"tpot": {
"mean": 0.14093466935703095,
"p50": 0.11608679958107468,
"p90": 0.24292321061978972,
"p99": 0.7622579889720426
},
"ttft": {
"mean": 0.4989603034128753,
"p50": 0.11800815758794723,
"p90": 1.394773550999063,
"p99": 5.683406566684342
}
}
],
"trialwise_distribution_bias": [
{
"e2e": {
"mean": 0.08560139205995335,
"p50": 0.11789837366536833,
"p90": 0.040257153080803736,
"p99": 0.05032404535106296
},
"tpot": {
"mean": 0.09063249747878419,
"p50": 0.10419652159558444,
"p90": 0.04737458218848663,
"p99": 0.03642414837746145
},
"ttft": {
"mean": 0.03099819883665655,
"p50": -0.0036305394992732502,
"p90": -0.03616773291859882,
"p99": -0.0027174231630140012
}
},
{
"e2e": {
"mean": 0.09717193562089739,
"p50": 0.12188146428139615,
"p90": 0.05076425771763047,
"p99": 0.05926112994918994
},
"tpot": {
"mean": 0.10315970439292232,
"p50": 0.11583781275404642,
"p90": 0.06417341590032893,
"p99": 0.037663262642925396
},
"ttft": {
"mean": 0.031678555957153555,
"p50": -0.001606874761951341,
"p90": -0.03560325750951616,
"p99": -0.0013960775328830752
}
}
],
"trialwise_distribution_bias_summary": {
"e2e": {
"mean": {
"max": 0.09717193562089739,
"mean": 0.09138666384042538,
"min": 0.08560139205995335
},
"p50": {
"max": 0.12188146428139615,
"mean": 0.11988991897338225,
"min": 0.11789837366536833
},
"p90": {
"max": 0.05076425771763047,
"mean": 0.0455107053992171,
"min": 0.040257153080803736
},
"p99": {
"max": 0.05926112994918994,
"mean": 0.05479258765012645,
"min": 0.05032404535106296
}
},
"tpot": {
"mean": {
"max": 0.10315970439292232,
"mean": 0.09689610093585325,
"min": 0.09063249747878419
},
"p50": {
"max": 0.11583781275404642,
"mean": 0.11001716717481544,
"min": 0.10419652159558444
},
"p90": {
"max": 0.06417341590032893,
"mean": 0.05577399904440778,
"min": 0.04737458218848663
},
"p99": {
"max": 0.037663262642925396,
"mean": 0.037043705510193425,
"min": 0.03642414837746145
}
},
"ttft": {
"mean": {
"max": 0.031678555957153555,
"mean": 0.03133837739690505,
"min": 0.03099819883665655
},
"p50": {
"max": -0.001606874761951341,
"mean": -0.002618707130612296,
"min": -0.0036305394992732502
},
"p90": {
"max": -0.03560325750951616,
"mean": -0.03588549521405749,
"min": -0.03616773291859882
},
"p99": {
"max": -0.0013960775328830752,
"mean": -0.0020567503479485385,
"min": -0.0027174231630140012
}
}
},
"waiting_p99_ms": 1296.3971075780032
}
}
},
"decision": "reject_global_scale_and_fit_shape_conditioned_curve",
"gates": {
"subcritical_mean_ttft_abs_bias_improves_3pp": false,
"subcritical_ttft_e2e_abs_bias_not_worse": false,
"subcritical_waiting_below_1s": true
},
"schema": "frontier-tp2-prefill-serving-replay-verdict.v1"
}

View File

@@ -0,0 +1,32 @@
{
"contract": {
"moe_mapping": "standalone moe_grouped_gemm and serving MoE both include expert prepare/finalize plus expert GEMMs",
"real_anchor_ms": 410.0,
"serving_sample": "longest execute window on critical TP rank",
"simulator_sample": "mean of first nine q8192/ctx0 single-request chunks"
},
"counterfactual": {
"moe_only_counterfactual_bias": -0.02523382099105694,
"moe_only_counterfactual_total_ms": 399.65413339366665,
"moe_residual_fraction": 0.8075231403622317,
"moe_residual_ms": 43.40535640199994,
"moe_scale": 1.2536568227334743,
"total_residual_ms": 53.751223008333284
},
"decision": "inject_tp2_prefill_moe_and_replay",
"gates": {
"moe_explains_at_least_70pct": true,
"moe_only_counterfactual_within_5pct": true,
"serving_reproduces_anchor_within_5pct": true
},
"measurements_ms": {
"real_anchor": 410.0,
"serving_execute_wall": 408.19065,
"serving_moe": 214.52378299999995,
"serving_non_moe": 193.66686700000005,
"simulator_moe_grouped_gemm": 171.118426598,
"simulator_non_moe": 185.1303503936667,
"simulator_total": 356.2487769916667
},
"schema": "frontier-tp2-prefill-serving-smoke-verdict.v1"
}

View File

@@ -0,0 +1,463 @@
{
"contract": {
"component_time": "sum of CUDA kernel duration within selected window",
"critical_path": "rank with largest selected execute wall",
"selection": "longest execute annotation per TP rank"
},
"critical_rank": {
"all_execute_windows": [
{
"duration_ms": 408.19065,
"name": "execute_context_1(8192)_generation_0(0)"
},
{
"duration_ms": 5.17557,
"name": "execute_context_0(0)_generation_1(1)"
}
],
"components_ms": {
"attention": 99.67212900000001,
"collective": 27.829565000000006,
"linear_norm_rope": 57.39125199999998,
"moe": 214.52378299999995,
"other": 1.9318720000000007,
"router": 0.7856679999999999
},
"execute_annotation_histogram": {
"execute_context_0(0)_generation_1(1)": 1,
"execute_context_1(8192)_generation_0(0)": 1
},
"execute_wall_ms": 408.19065,
"gpu_kernel_busy_ms": 402.13426899999996,
"kernel_rows": [
{
"duration_ms": 122.02009999999999,
"name": "void fused_moe::run_global<fused_moe::Fused_Moe_Kernel_sm80<cutlass::bfloat16_t, cutlass::bfloat16_t, cutlass::bfloat16_t, 32, 128, 64, 3, (fused_moe::Activation_Type)3> >(fused_moe::Fused_Moe_Kernel_sm80<cutlass::bfloat16_t, cutlass::bfloat16_t, cutlass::bfloat16_t, 32, 128, 64, 3, (fused_moe::Activation_Type)3>::Params)"
},
{
"duration_ms": 97.15557600000001,
"name": "void cutlass::device_kernel<flash::enable_sm90_or_later<flash::FlashAttnFwdSm90<flash::CollectiveMainloopFwdSm90<2, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1> >, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, 128, cutlass::bfloat16_t, float, cutlass::arch::Sm90, true, false, false, true, true, false, false, true, true, true, false, false, cutlass::bfloat16_t, 8>, flash::CollectiveEpilogueFwd<cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1> >, cutlass::bfloat16_t, cutlass::arch::Sm90, 256, true, true, false, false, 8>, flash::VarlenDynamicPersistentTileScheduler<128, 128, 256, 128, false, true, true, true, false, true> > > >(flash::enable_sm90_or_later<flash::FlashAttnFwdSm90<flash::CollectiveMainloopFwdSm90<2, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1> >, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, 128, cutlass::bfloat16_t, float, cutlass::arch::Sm90, true, false, false, true, true, false, false, true, true, true, false, false, cutlass::bfloat16_t, 8>, flash::CollectiveEpilogueFwd<cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1> >, cutlass::bfloat16_t, cutlass::arch::Sm90, 256, true, true, false, false, 8>, flash::VarlenDynamicPersistentTileScheduler<128, 128, 256, 128, false, true, true, true, false, true> > >::Params)"
},
{
"duration_ms": 80.135489,
"name": "_ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalINS1_17GroupProblemShapeIN4cute5tupleIJlllEEEEENS1_10collective13CollectiveMmaINS1_39MainloopSm90ArrayTmaGmmaWarpSpecializedILi12ENS6_IJNS5_1CILi1EEESD_SD_EEENS1_43KernelPtrArrayTmaWarpSpecializedCooperativeEEENS6_IJNSC_ILi128EEENSC_ILi16EEENSC_ILi64EEEEEENS_10bfloat16_tEPNS6_IJlSD_NSC_ILi0EEEEEESL_SO_NS5_8TiledMMAINS5_8MMA_AtomIJNS5_4SM904GMMA27MMA_64x16x16_F32BF16BF16_SSILNSS_5MajorE0ELSU_0ELNSS_7ScaleInE1ELSV_1EEEEEENS5_6LayoutINS6_IJNSC_ILi2EEESD_SD_EEENS6_IJSD_SM_SM_EEEEENS6_IJNS5_10UnderscoreES13_S13_EEEEENS5_13SM90_TMA_LOADENS5_14ComposedLayoutINS5_7SwizzleILi3ELi4ELi3EEENS5_18smem_ptr_flag_bitsILi16EEENSY_INS6_IJNSC_ILi8EEESJ_EEENS6_IJSJ_SD_EEEEEEEvNS5_8identityES16_S1G_vS1H_EENS_8epilogue10collective18CollectiveEpilogueINS1J_30Sm90PtrArrayTmaWarpSpecializedILi1ELi1ELi8ELb0ELb0ELi2EEEJSK_NS6_IJSH_SI_EEEvPNS6_IJSD_lSM_EEEvS1Q_NS1J_6fusion15FusionCallbacksIS1N_NS1R_37ScaledAccPerRowBiasPerColScaleScatterINS_6layout11ColumnMajorESL_fSL_ffLi8ELi8ELNS_15FloatRoundStyleE2EEESK_S1O_JNS17_IS19_S1B_NSY_INS6_IJSJ_S1C_EEENS6_IJSD_SJ_EEEEEEENS5_17SM90_U16x8_STSM_TEEEES16_S21_NS5_17SM75_U16x8_LDSM_TENS5_14SM90_TMA_STOREES21_S22_NS5_9Copy_AtomIJNS5_17SM90_U32x4_STSM_NENS_6half_tEEEEvEEEvvEEEEvNT_6ParamsE"
},
{
"duration_ms": 31.169920999999995,
"name": "nvjet_tst_320x128_64x3_1x2_h_bz_coopB_TNT"
},
{
"duration_ms": 27.829565000000006,
"name": "void flashinfer::trtllm_allreduce_fusion::allreduce_fusion_kernel_oneshot_lamport<(flashinfer::trtllm_allreduce_fusion::AllReduceFusionPattern)1, __nv_bfloat16, 2, true, true>(flashinfer::trtllm_allreduce_fusion::AllReduceFusionParams<__nv_bfloat16>)"
},
{
"duration_ms": 24.150314000000005,
"name": "nvjet_tst_128x192_64x5_2x1_v_bz_coopB_TNN"
},
{
"duration_ms": 9.48525,
"name": "void tensorrt_llm::kernels::cutlass_kernels::expandInputRowsKernel<__nv_bfloat16, __nv_bfloat16, (tensorrt_llm::kernels::cutlass_kernels::TmaWarpSpecializedGroupedGemmInput::FpXBlockScalingType)2, false>(__nv_bfloat16 const*, __nv_bfloat16*, float const*, float*, int const*, long, long, long, float const*, bool, long const*, unsigned char*, unsigned char const*, bool, long, __nv_bfloat16 const*)"
},
{
"duration_ms": 2.5150810000000003,
"name": "void vllm::reshape_and_cache_flash_kernel<__nv_bfloat16, __nv_bfloat16, (vllm::Fp8KVCacheDataType)0>(__nv_bfloat16 const*, __nv_bfloat16 const*, __nv_bfloat16*, __nv_bfloat16*, long const*, long, long, long, long, long, int, int, int, float const*, float const*, int)"
},
{
"duration_ms": 1.9769679999999998,
"name": "nvjet_tst_128x128_64x6_1x2_h_bz_TNT"
},
{
"duration_ms": 1.356286,
"name": "void tensorrt_llm::kernels::cutlass_kernels::blockExpertPrefixSumKernel<1024>(int const*, int*, int*, long, long, int)"
},
{
"duration_ms": 1.1989450000000001,
"name": "triton_poi_fused_1"
},
{
"duration_ms": 1.1752029999999996,
"name": "void tensorrt_llm::kernels::cutlass_kernels::computeStridesTmaWarpSpecializedKernel<__nv_bfloat16, __nv_bfloat16, __nv_bfloat16, __nv_bfloat16>(long const*, tensorrt_llm::kernels::cutlass_kernels::TmaWarpSpecializedGroupedGemmInput, tensorrt_llm::kernels::cutlass_kernels::TmaWarpSpecializedGroupedGemmInput, long, long, long, long, long, long, long, __nv_bfloat16 const*, __nv_bfloat16 const*, __nv_bfloat16 const*, __nv_bfloat16 const*, float const*, float const*, unsigned char const*, unsigned char const*, tensorrt_llm::kernels::cutlass_kernels::QuantParams, __nv_bfloat16 const*, __nv_bfloat16 const*, __nv_bfloat16*, __nv_bfloat16*, float const*, int const*)"
},
{
"duration_ms": 0.7856679999999999,
"name": "void vllm::moe::topkGating<8, 128, 4, 16, 32, int, __nv_bfloat16, (vllm::moe::ScoringFunc)0>(__nv_bfloat16 const*, bool const*, float*, int, int*, int*, int, int, int, bool, float const*)"
},
{
"duration_ms": 0.6054410000000001,
"name": "triton_red_fused_0"
},
{
"duration_ms": 0.264768,
"name": "tensorrt_llm::kernels::cutlass_kernels::mergeExpertPrefixSumKernel(int const*, int const*, int const*, int*, int*, int*, int)"
},
{
"duration_ms": 0.09404900000000001,
"name": "nvjet_tst_512x8_64x3_2x1_v_bz_TNT"
},
{
"duration_ms": 0.086687,
"name": "void tensorrt_llm::kernels::cutlass_kernels::globalExpertPrefixSumKernel<1024>(int const*, int*, long*, long, long)"
},
{
"duration_ms": 0.034815,
"name": "void at::native::vectorized_elementwise_kernel<4, at::native::FillFunctor<int>, std::array<char*, 1ul> >(int, at::native::FillFunctor<int>, std::array<char*, 1ul>)"
},
{
"duration_ms": 0.025472,
"name": "triton_poi_fused_2"
},
{
"duration_ms": 0.022335999999999998,
"name": "triton_poi_fused_0"
},
{
"duration_ms": 0.015392,
"name": "_compute_slot_mapping_kernel"
},
{
"duration_ms": 0.013439999999999999,
"name": "triton_red_fused_1"
},
{
"duration_ms": 0.004256,
"name": "void at::native::unrolled_elementwise_kernel<at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#4}::operator()() const::{lambda(long)#1}, std::array<char*, 2ul>, 4, TrivialOffsetCalculator<1, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithCast<1>, at::native::memory::StoreWithCast<1> >(int, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#4}::operator()() const::{lambda(long)#1}, std::array<char*, 2ul>, TrivialOffsetCalculator<1, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithCast<1>, at::native::memory::StoreWithCast<1>)"
},
{
"duration_ms": 0.0041600000000000005,
"name": "void at::native::index_elementwise_kernel<128, 4, at::native::gpu_index_kernel<at::native::index_kernel_impl<at::native::OpaqueType<4> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>, at::native::index_kernel_impl<at::native::OpaqueType<4> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1} const&, bool)::{lambda(int)#1}>(long, at::native::gpu_index_kernel<at::native::index_kernel_impl<at::native::OpaqueType<4> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>, at::native::index_kernel_impl<at::native::OpaqueType<4> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1} const&, bool)::{lambda(int)#1})"
},
{
"duration_ms": 0.003072,
"name": "void at::native::index_elementwise_kernel<128, 4, at::native::gpu_index_kernel<at::native::index_kernel_impl<at::native::OpaqueType<2> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>, at::native::index_kernel_impl<at::native::OpaqueType<2> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1} const&, bool)::{lambda(int)#1}>(long, at::native::gpu_index_kernel<at::native::index_kernel_impl<at::native::OpaqueType<2> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>, at::native::index_kernel_impl<at::native::OpaqueType<2> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1} const&, bool)::{lambda(int)#1})"
},
{
"duration_ms": 0.001472,
"name": "void at::native::vectorized_elementwise_kernel<2, at::native::CUDAFunctor_add<long>, std::array<char*, 3ul> >(int, at::native::CUDAFunctor_add<long>, std::array<char*, 3ul>)"
},
{
"duration_ms": 0.001472,
"name": "void flash::prepare_varlen_num_blocks_kernel<1, false>(int, int, int, int const*, int const*, int const*, int const*, int const*, int const*, int, int, int, int, int, cutlass::FastDivmod, cutlass::FastDivmod, int*, int*, int*, int*, int*, bool, bool, bool, int)"
},
{
"duration_ms": 0.001184,
"name": "void at::native::unrolled_elementwise_kernel<at::native::FillFunctor<int>, std::array<char*, 1ul>, 4, TrivialOffsetCalculator<0, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithoutCast, at::native::memory::StoreWithoutCast>(int, at::native::FillFunctor<int>, std::array<char*, 1ul>, TrivialOffsetCalculator<0, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithoutCast, at::native::memory::StoreWithoutCast)"
},
{
"duration_ms": 0.0011200000000000001,
"name": "void at::native::vectorized_elementwise_kernel<4, at::native::CUDAFunctor_add<int>, std::array<char*, 3ul> >(int, at::native::CUDAFunctor_add<int>, std::array<char*, 3ul>)"
},
{
"duration_ms": 0.000767,
"name": "void at::native::unrolled_elementwise_kernel<at::native::CUDAFunctorOnSelf_add<int>, std::array<char*, 2ul>, 4, TrivialOffsetCalculator<1, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithoutCast, at::native::memory::StoreWithoutCast>(int, at::native::CUDAFunctorOnSelf_add<int>, std::array<char*, 2ul>, TrivialOffsetCalculator<1, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithoutCast, at::native::memory::StoreWithoutCast)"
}
],
"non_kernel_gap_ms": 6.056381000000044,
"selected_execute_annotation": "execute_context_1(8192)_generation_0(0)",
"trace": "runs/frontier-tp2-prefill-serving-v0/fleet-artifacts/tp2-prefill-serving-smoke-20260723-v1-20260723T082454699036Z/artifacts/runs/frontier-tp2-prefill-serving-v0/remote-outputs/tp2-smoke-r1/traces/profile/dp0_pp0_tp0_dcp0_ep0_rank0.1784795312743204872.pt.trace.json.gz"
},
"ranks": [
{
"all_execute_windows": [
{
"duration_ms": 408.19065,
"name": "execute_context_1(8192)_generation_0(0)"
},
{
"duration_ms": 5.17557,
"name": "execute_context_0(0)_generation_1(1)"
}
],
"components_ms": {
"attention": 99.67212900000001,
"collective": 27.829565000000006,
"linear_norm_rope": 57.39125199999998,
"moe": 214.52378299999995,
"other": 1.9318720000000007,
"router": 0.7856679999999999
},
"execute_annotation_histogram": {
"execute_context_0(0)_generation_1(1)": 1,
"execute_context_1(8192)_generation_0(0)": 1
},
"execute_wall_ms": 408.19065,
"gpu_kernel_busy_ms": 402.13426899999996,
"kernel_rows": [
{
"duration_ms": 122.02009999999999,
"name": "void fused_moe::run_global<fused_moe::Fused_Moe_Kernel_sm80<cutlass::bfloat16_t, cutlass::bfloat16_t, cutlass::bfloat16_t, 32, 128, 64, 3, (fused_moe::Activation_Type)3> >(fused_moe::Fused_Moe_Kernel_sm80<cutlass::bfloat16_t, cutlass::bfloat16_t, cutlass::bfloat16_t, 32, 128, 64, 3, (fused_moe::Activation_Type)3>::Params)"
},
{
"duration_ms": 97.15557600000001,
"name": "void cutlass::device_kernel<flash::enable_sm90_or_later<flash::FlashAttnFwdSm90<flash::CollectiveMainloopFwdSm90<2, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1> >, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, 128, cutlass::bfloat16_t, float, cutlass::arch::Sm90, true, false, false, true, true, false, false, true, true, true, false, false, cutlass::bfloat16_t, 8>, flash::CollectiveEpilogueFwd<cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1> >, cutlass::bfloat16_t, cutlass::arch::Sm90, 256, true, true, false, false, 8>, flash::VarlenDynamicPersistentTileScheduler<128, 128, 256, 128, false, true, true, true, false, true> > > >(flash::enable_sm90_or_later<flash::FlashAttnFwdSm90<flash::CollectiveMainloopFwdSm90<2, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1> >, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, 128, cutlass::bfloat16_t, float, cutlass::arch::Sm90, true, false, false, true, true, false, false, true, true, true, false, false, cutlass::bfloat16_t, 8>, flash::CollectiveEpilogueFwd<cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1> >, cutlass::bfloat16_t, cutlass::arch::Sm90, 256, true, true, false, false, 8>, flash::VarlenDynamicPersistentTileScheduler<128, 128, 256, 128, false, true, true, true, false, true> > >::Params)"
},
{
"duration_ms": 80.135489,
"name": "_ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalINS1_17GroupProblemShapeIN4cute5tupleIJlllEEEEENS1_10collective13CollectiveMmaINS1_39MainloopSm90ArrayTmaGmmaWarpSpecializedILi12ENS6_IJNS5_1CILi1EEESD_SD_EEENS1_43KernelPtrArrayTmaWarpSpecializedCooperativeEEENS6_IJNSC_ILi128EEENSC_ILi16EEENSC_ILi64EEEEEENS_10bfloat16_tEPNS6_IJlSD_NSC_ILi0EEEEEESL_SO_NS5_8TiledMMAINS5_8MMA_AtomIJNS5_4SM904GMMA27MMA_64x16x16_F32BF16BF16_SSILNSS_5MajorE0ELSU_0ELNSS_7ScaleInE1ELSV_1EEEEEENS5_6LayoutINS6_IJNSC_ILi2EEESD_SD_EEENS6_IJSD_SM_SM_EEEEENS6_IJNS5_10UnderscoreES13_S13_EEEEENS5_13SM90_TMA_LOADENS5_14ComposedLayoutINS5_7SwizzleILi3ELi4ELi3EEENS5_18smem_ptr_flag_bitsILi16EEENSY_INS6_IJNSC_ILi8EEESJ_EEENS6_IJSJ_SD_EEEEEEEvNS5_8identityES16_S1G_vS1H_EENS_8epilogue10collective18CollectiveEpilogueINS1J_30Sm90PtrArrayTmaWarpSpecializedILi1ELi1ELi8ELb0ELb0ELi2EEEJSK_NS6_IJSH_SI_EEEvPNS6_IJSD_lSM_EEEvS1Q_NS1J_6fusion15FusionCallbacksIS1N_NS1R_37ScaledAccPerRowBiasPerColScaleScatterINS_6layout11ColumnMajorESL_fSL_ffLi8ELi8ELNS_15FloatRoundStyleE2EEESK_S1O_JNS17_IS19_S1B_NSY_INS6_IJSJ_S1C_EEENS6_IJSD_SJ_EEEEEEENS5_17SM90_U16x8_STSM_TEEEES16_S21_NS5_17SM75_U16x8_LDSM_TENS5_14SM90_TMA_STOREES21_S22_NS5_9Copy_AtomIJNS5_17SM90_U32x4_STSM_NENS_6half_tEEEEvEEEvvEEEEvNT_6ParamsE"
},
{
"duration_ms": 31.169920999999995,
"name": "nvjet_tst_320x128_64x3_1x2_h_bz_coopB_TNT"
},
{
"duration_ms": 27.829565000000006,
"name": "void flashinfer::trtllm_allreduce_fusion::allreduce_fusion_kernel_oneshot_lamport<(flashinfer::trtllm_allreduce_fusion::AllReduceFusionPattern)1, __nv_bfloat16, 2, true, true>(flashinfer::trtllm_allreduce_fusion::AllReduceFusionParams<__nv_bfloat16>)"
},
{
"duration_ms": 24.150314000000005,
"name": "nvjet_tst_128x192_64x5_2x1_v_bz_coopB_TNN"
},
{
"duration_ms": 9.48525,
"name": "void tensorrt_llm::kernels::cutlass_kernels::expandInputRowsKernel<__nv_bfloat16, __nv_bfloat16, (tensorrt_llm::kernels::cutlass_kernels::TmaWarpSpecializedGroupedGemmInput::FpXBlockScalingType)2, false>(__nv_bfloat16 const*, __nv_bfloat16*, float const*, float*, int const*, long, long, long, float const*, bool, long const*, unsigned char*, unsigned char const*, bool, long, __nv_bfloat16 const*)"
},
{
"duration_ms": 2.5150810000000003,
"name": "void vllm::reshape_and_cache_flash_kernel<__nv_bfloat16, __nv_bfloat16, (vllm::Fp8KVCacheDataType)0>(__nv_bfloat16 const*, __nv_bfloat16 const*, __nv_bfloat16*, __nv_bfloat16*, long const*, long, long, long, long, long, int, int, int, float const*, float const*, int)"
},
{
"duration_ms": 1.9769679999999998,
"name": "nvjet_tst_128x128_64x6_1x2_h_bz_TNT"
},
{
"duration_ms": 1.356286,
"name": "void tensorrt_llm::kernels::cutlass_kernels::blockExpertPrefixSumKernel<1024>(int const*, int*, int*, long, long, int)"
},
{
"duration_ms": 1.1989450000000001,
"name": "triton_poi_fused_1"
},
{
"duration_ms": 1.1752029999999996,
"name": "void tensorrt_llm::kernels::cutlass_kernels::computeStridesTmaWarpSpecializedKernel<__nv_bfloat16, __nv_bfloat16, __nv_bfloat16, __nv_bfloat16>(long const*, tensorrt_llm::kernels::cutlass_kernels::TmaWarpSpecializedGroupedGemmInput, tensorrt_llm::kernels::cutlass_kernels::TmaWarpSpecializedGroupedGemmInput, long, long, long, long, long, long, long, __nv_bfloat16 const*, __nv_bfloat16 const*, __nv_bfloat16 const*, __nv_bfloat16 const*, float const*, float const*, unsigned char const*, unsigned char const*, tensorrt_llm::kernels::cutlass_kernels::QuantParams, __nv_bfloat16 const*, __nv_bfloat16 const*, __nv_bfloat16*, __nv_bfloat16*, float const*, int const*)"
},
{
"duration_ms": 0.7856679999999999,
"name": "void vllm::moe::topkGating<8, 128, 4, 16, 32, int, __nv_bfloat16, (vllm::moe::ScoringFunc)0>(__nv_bfloat16 const*, bool const*, float*, int, int*, int*, int, int, int, bool, float const*)"
},
{
"duration_ms": 0.6054410000000001,
"name": "triton_red_fused_0"
},
{
"duration_ms": 0.264768,
"name": "tensorrt_llm::kernels::cutlass_kernels::mergeExpertPrefixSumKernel(int const*, int const*, int const*, int*, int*, int*, int)"
},
{
"duration_ms": 0.09404900000000001,
"name": "nvjet_tst_512x8_64x3_2x1_v_bz_TNT"
},
{
"duration_ms": 0.086687,
"name": "void tensorrt_llm::kernels::cutlass_kernels::globalExpertPrefixSumKernel<1024>(int const*, int*, long*, long, long)"
},
{
"duration_ms": 0.034815,
"name": "void at::native::vectorized_elementwise_kernel<4, at::native::FillFunctor<int>, std::array<char*, 1ul> >(int, at::native::FillFunctor<int>, std::array<char*, 1ul>)"
},
{
"duration_ms": 0.025472,
"name": "triton_poi_fused_2"
},
{
"duration_ms": 0.022335999999999998,
"name": "triton_poi_fused_0"
},
{
"duration_ms": 0.015392,
"name": "_compute_slot_mapping_kernel"
},
{
"duration_ms": 0.013439999999999999,
"name": "triton_red_fused_1"
},
{
"duration_ms": 0.004256,
"name": "void at::native::unrolled_elementwise_kernel<at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#4}::operator()() const::{lambda(long)#1}, std::array<char*, 2ul>, 4, TrivialOffsetCalculator<1, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithCast<1>, at::native::memory::StoreWithCast<1> >(int, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#4}::operator()() const::{lambda(long)#1}, std::array<char*, 2ul>, TrivialOffsetCalculator<1, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithCast<1>, at::native::memory::StoreWithCast<1>)"
},
{
"duration_ms": 0.0041600000000000005,
"name": "void at::native::index_elementwise_kernel<128, 4, at::native::gpu_index_kernel<at::native::index_kernel_impl<at::native::OpaqueType<4> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>, at::native::index_kernel_impl<at::native::OpaqueType<4> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1} const&, bool)::{lambda(int)#1}>(long, at::native::gpu_index_kernel<at::native::index_kernel_impl<at::native::OpaqueType<4> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>, at::native::index_kernel_impl<at::native::OpaqueType<4> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1} const&, bool)::{lambda(int)#1})"
},
{
"duration_ms": 0.003072,
"name": "void at::native::index_elementwise_kernel<128, 4, at::native::gpu_index_kernel<at::native::index_kernel_impl<at::native::OpaqueType<2> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>, at::native::index_kernel_impl<at::native::OpaqueType<2> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1} const&, bool)::{lambda(int)#1}>(long, at::native::gpu_index_kernel<at::native::index_kernel_impl<at::native::OpaqueType<2> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>, at::native::index_kernel_impl<at::native::OpaqueType<2> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1} const&, bool)::{lambda(int)#1})"
},
{
"duration_ms": 0.001472,
"name": "void at::native::vectorized_elementwise_kernel<2, at::native::CUDAFunctor_add<long>, std::array<char*, 3ul> >(int, at::native::CUDAFunctor_add<long>, std::array<char*, 3ul>)"
},
{
"duration_ms": 0.001472,
"name": "void flash::prepare_varlen_num_blocks_kernel<1, false>(int, int, int, int const*, int const*, int const*, int const*, int const*, int const*, int, int, int, int, int, cutlass::FastDivmod, cutlass::FastDivmod, int*, int*, int*, int*, int*, bool, bool, bool, int)"
},
{
"duration_ms": 0.001184,
"name": "void at::native::unrolled_elementwise_kernel<at::native::FillFunctor<int>, std::array<char*, 1ul>, 4, TrivialOffsetCalculator<0, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithoutCast, at::native::memory::StoreWithoutCast>(int, at::native::FillFunctor<int>, std::array<char*, 1ul>, TrivialOffsetCalculator<0, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithoutCast, at::native::memory::StoreWithoutCast)"
},
{
"duration_ms": 0.0011200000000000001,
"name": "void at::native::vectorized_elementwise_kernel<4, at::native::CUDAFunctor_add<int>, std::array<char*, 3ul> >(int, at::native::CUDAFunctor_add<int>, std::array<char*, 3ul>)"
},
{
"duration_ms": 0.000767,
"name": "void at::native::unrolled_elementwise_kernel<at::native::CUDAFunctorOnSelf_add<int>, std::array<char*, 2ul>, 4, TrivialOffsetCalculator<1, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithoutCast, at::native::memory::StoreWithoutCast>(int, at::native::CUDAFunctorOnSelf_add<int>, std::array<char*, 2ul>, TrivialOffsetCalculator<1, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithoutCast, at::native::memory::StoreWithoutCast)"
}
],
"non_kernel_gap_ms": 6.056381000000044,
"selected_execute_annotation": "execute_context_1(8192)_generation_0(0)",
"trace": "runs/frontier-tp2-prefill-serving-v0/fleet-artifacts/tp2-prefill-serving-smoke-20260723-v1-20260723T082454699036Z/artifacts/runs/frontier-tp2-prefill-serving-v0/remote-outputs/tp2-smoke-r1/traces/profile/dp0_pp0_tp0_dcp0_ep0_rank0.1784795312743204872.pt.trace.json.gz"
},
{
"all_execute_windows": [
{
"duration_ms": 407.337292,
"name": "execute_context_1(8192)_generation_0(0)"
},
{
"duration_ms": 5.175196,
"name": "execute_context_0(0)_generation_1(1)"
}
],
"components_ms": {
"attention": 99.75531199999998,
"collective": 25.666238000000003,
"linear_norm_rope": 57.43388800000001,
"moe": 214.55563300000009,
"other": 1.9427889999999999,
"router": 0.7898920000000001
},
"execute_annotation_histogram": {
"execute_context_0(0)_generation_1(1)": 1,
"execute_context_1(8192)_generation_0(0)": 1
},
"execute_wall_ms": 407.337292,
"gpu_kernel_busy_ms": 400.14375200000006,
"kernel_rows": [
{
"duration_ms": 122.088157,
"name": "void fused_moe::run_global<fused_moe::Fused_Moe_Kernel_sm80<cutlass::bfloat16_t, cutlass::bfloat16_t, cutlass::bfloat16_t, 32, 128, 64, 3, (fused_moe::Activation_Type)3> >(fused_moe::Fused_Moe_Kernel_sm80<cutlass::bfloat16_t, cutlass::bfloat16_t, cutlass::bfloat16_t, 32, 128, 64, 3, (fused_moe::Activation_Type)3>::Params)"
},
{
"duration_ms": 97.24454899999998,
"name": "void cutlass::device_kernel<flash::enable_sm90_or_later<flash::FlashAttnFwdSm90<flash::CollectiveMainloopFwdSm90<2, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1> >, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, 128, cutlass::bfloat16_t, float, cutlass::arch::Sm90, true, false, false, true, true, false, false, true, true, true, false, false, cutlass::bfloat16_t, 8>, flash::CollectiveEpilogueFwd<cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1> >, cutlass::bfloat16_t, cutlass::arch::Sm90, 256, true, true, false, false, 8>, flash::VarlenDynamicPersistentTileScheduler<128, 128, 256, 128, false, true, true, true, false, true> > > >(flash::enable_sm90_or_later<flash::FlashAttnFwdSm90<flash::CollectiveMainloopFwdSm90<2, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1> >, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, 128, cutlass::bfloat16_t, float, cutlass::arch::Sm90, true, false, false, true, true, false, false, true, true, true, false, false, cutlass::bfloat16_t, 8>, flash::CollectiveEpilogueFwd<cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1> >, cutlass::bfloat16_t, cutlass::arch::Sm90, 256, true, true, false, false, 8>, flash::VarlenDynamicPersistentTileScheduler<128, 128, 256, 128, false, true, true, true, false, true> > >::Params)"
},
{
"duration_ms": 80.12219199999998,
"name": "_ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalINS1_17GroupProblemShapeIN4cute5tupleIJlllEEEEENS1_10collective13CollectiveMmaINS1_39MainloopSm90ArrayTmaGmmaWarpSpecializedILi12ENS6_IJNS5_1CILi1EEESD_SD_EEENS1_43KernelPtrArrayTmaWarpSpecializedCooperativeEEENS6_IJNSC_ILi128EEENSC_ILi16EEENSC_ILi64EEEEEENS_10bfloat16_tEPNS6_IJlSD_NSC_ILi0EEEEEESL_SO_NS5_8TiledMMAINS5_8MMA_AtomIJNS5_4SM904GMMA27MMA_64x16x16_F32BF16BF16_SSILNSS_5MajorE0ELSU_0ELNSS_7ScaleInE1ELSV_1EEEEEENS5_6LayoutINS6_IJNSC_ILi2EEESD_SD_EEENS6_IJSD_SM_SM_EEEEENS6_IJNS5_10UnderscoreES13_S13_EEEEENS5_13SM90_TMA_LOADENS5_14ComposedLayoutINS5_7SwizzleILi3ELi4ELi3EEENS5_18smem_ptr_flag_bitsILi16EEENSY_INS6_IJNSC_ILi8EEESJ_EEENS6_IJSJ_SD_EEEEEEEvNS5_8identityES16_S1G_vS1H_EENS_8epilogue10collective18CollectiveEpilogueINS1J_30Sm90PtrArrayTmaWarpSpecializedILi1ELi1ELi8ELb0ELb0ELi2EEEJSK_NS6_IJSH_SI_EEEvPNS6_IJSD_lSM_EEEvS1Q_NS1J_6fusion15FusionCallbacksIS1N_NS1R_37ScaledAccPerRowBiasPerColScaleScatterINS_6layout11ColumnMajorESL_fSL_ffLi8ELi8ELNS_15FloatRoundStyleE2EEESK_S1O_JNS17_IS19_S1B_NSY_INS6_IJSJ_S1C_EEENS6_IJSD_SJ_EEEEEEENS5_17SM90_U16x8_STSM_TEEEES16_S21_NS5_17SM75_U16x8_LDSM_TENS5_14SM90_TMA_STOREES21_S22_NS5_9Copy_AtomIJNS5_17SM90_U32x4_STSM_NENS_6half_tEEEEvEEEvvEEEEvNT_6ParamsE"
},
{
"duration_ms": 31.199706,
"name": "nvjet_tst_320x128_64x3_1x2_h_bz_coopB_TNT"
},
{
"duration_ms": 25.666238000000003,
"name": "void flashinfer::trtllm_allreduce_fusion::allreduce_fusion_kernel_oneshot_lamport<(flashinfer::trtllm_allreduce_fusion::AllReduceFusionPattern)1, __nv_bfloat16, 2, true, true>(flashinfer::trtllm_allreduce_fusion::AllReduceFusionParams<__nv_bfloat16>)"
},
{
"duration_ms": 24.162201000000007,
"name": "nvjet_tst_128x192_64x5_2x1_v_bz_coopB_TNN"
},
{
"duration_ms": 9.475503000000002,
"name": "void tensorrt_llm::kernels::cutlass_kernels::expandInputRowsKernel<__nv_bfloat16, __nv_bfloat16, (tensorrt_llm::kernels::cutlass_kernels::TmaWarpSpecializedGroupedGemmInput::FpXBlockScalingType)2, false>(__nv_bfloat16 const*, __nv_bfloat16*, float const*, float*, int const*, long, long, long, float const*, bool, long const*, unsigned char*, unsigned char const*, bool, long, __nv_bfloat16 const*)"
},
{
"duration_ms": 2.5092910000000006,
"name": "void vllm::reshape_and_cache_flash_kernel<__nv_bfloat16, __nv_bfloat16, (vllm::Fp8KVCacheDataType)0>(__nv_bfloat16 const*, __nv_bfloat16 const*, __nv_bfloat16*, __nv_bfloat16*, long const*, long, long, long, long, long, int, int, int, float const*, float const*, int)"
},
{
"duration_ms": 1.9775490000000004,
"name": "nvjet_tst_128x128_64x6_1x2_h_bz_TNT"
},
{
"duration_ms": 1.3556549999999998,
"name": "void tensorrt_llm::kernels::cutlass_kernels::blockExpertPrefixSumKernel<1024>(int const*, int*, int*, long, long, int)"
},
{
"duration_ms": 1.2008649999999996,
"name": "triton_poi_fused_1"
},
{
"duration_ms": 1.164903,
"name": "void tensorrt_llm::kernels::cutlass_kernels::computeStridesTmaWarpSpecializedKernel<__nv_bfloat16, __nv_bfloat16, __nv_bfloat16, __nv_bfloat16>(long const*, tensorrt_llm::kernels::cutlass_kernels::TmaWarpSpecializedGroupedGemmInput, tensorrt_llm::kernels::cutlass_kernels::TmaWarpSpecializedGroupedGemmInput, long, long, long, long, long, long, long, __nv_bfloat16 const*, __nv_bfloat16 const*, __nv_bfloat16 const*, __nv_bfloat16 const*, float const*, float const*, unsigned char const*, unsigned char const*, tensorrt_llm::kernels::cutlass_kernels::QuantParams, __nv_bfloat16 const*, __nv_bfloat16 const*, __nv_bfloat16*, __nv_bfloat16*, float const*, int const*)"
},
{
"duration_ms": 0.7898920000000001,
"name": "void vllm::moe::topkGating<8, 128, 4, 16, 32, int, __nv_bfloat16, (vllm::moe::ScoringFunc)0>(__nv_bfloat16 const*, bool const*, float*, int, int*, int*, int, int, int, bool, float const*)"
},
{
"duration_ms": 0.6094719999999998,
"name": "triton_red_fused_0"
},
{
"duration_ms": 0.2641649999999999,
"name": "tensorrt_llm::kernels::cutlass_kernels::mergeExpertPrefixSumKernel(int const*, int const*, int const*, int*, int*, int*, int)"
},
{
"duration_ms": 0.094432,
"name": "nvjet_tst_512x8_64x3_2x1_v_bz_TNT"
},
{
"duration_ms": 0.08505799999999997,
"name": "void tensorrt_llm::kernels::cutlass_kernels::globalExpertPrefixSumKernel<1024>(int const*, int*, long*, long, long)"
},
{
"duration_ms": 0.034820000000000004,
"name": "void at::native::vectorized_elementwise_kernel<4, at::native::FillFunctor<int>, std::array<char*, 1ul> >(int, at::native::FillFunctor<int>, std::array<char*, 1ul>)"
},
{
"duration_ms": 0.027615999999999998,
"name": "triton_poi_fused_0"
},
{
"duration_ms": 0.025792000000000002,
"name": "triton_poi_fused_2"
},
{
"duration_ms": 0.015392,
"name": "_compute_slot_mapping_kernel"
},
{
"duration_ms": 0.012928,
"name": "triton_red_fused_1"
},
{
"duration_ms": 0.004096,
"name": "void at::native::unrolled_elementwise_kernel<at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#4}::operator()() const::{lambda(long)#1}, std::array<char*, 2ul>, 4, TrivialOffsetCalculator<1, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithCast<1>, at::native::memory::StoreWithCast<1> >(int, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#4}::operator()() const::{lambda(long)#1}, std::array<char*, 2ul>, TrivialOffsetCalculator<1, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithCast<1>, at::native::memory::StoreWithCast<1>)"
},
{
"duration_ms": 0.004,
"name": "void at::native::index_elementwise_kernel<128, 4, at::native::gpu_index_kernel<at::native::index_kernel_impl<at::native::OpaqueType<4> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>, at::native::index_kernel_impl<at::native::OpaqueType<4> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1} const&, bool)::{lambda(int)#1}>(long, at::native::gpu_index_kernel<at::native::index_kernel_impl<at::native::OpaqueType<4> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>, at::native::index_kernel_impl<at::native::OpaqueType<4> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1} const&, bool)::{lambda(int)#1})"
},
{
"duration_ms": 0.003328,
"name": "void at::native::index_elementwise_kernel<128, 4, at::native::gpu_index_kernel<at::native::index_kernel_impl<at::native::OpaqueType<2> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>, at::native::index_kernel_impl<at::native::OpaqueType<2> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1} const&, bool)::{lambda(int)#1}>(long, at::native::gpu_index_kernel<at::native::index_kernel_impl<at::native::OpaqueType<2> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>, at::native::index_kernel_impl<at::native::OpaqueType<2> >(at::TensorIteratorBase&, c10::ArrayRef<long>, c10::ArrayRef<long>)::{lambda(char*, char const*, long)#1} const&, bool)::{lambda(int)#1})"
},
{
"duration_ms": 0.001472,
"name": "void flash::prepare_varlen_num_blocks_kernel<1, false>(int, int, int, int const*, int const*, int const*, int const*, int const*, int const*, int, int, int, int, int, cutlass::FastDivmod, cutlass::FastDivmod, int*, int*, int*, int*, int*, bool, bool, bool, int)"
},
{
"duration_ms": 0.001312,
"name": "void at::native::vectorized_elementwise_kernel<2, at::native::CUDAFunctor_add<long>, std::array<char*, 3ul> >(int, at::native::CUDAFunctor_add<long>, std::array<char*, 3ul>)"
},
{
"duration_ms": 0.001184,
"name": "void at::native::vectorized_elementwise_kernel<4, at::native::CUDAFunctor_add<int>, std::array<char*, 3ul> >(int, at::native::CUDAFunctor_add<int>, std::array<char*, 3ul>)"
},
{
"duration_ms": 0.001184,
"name": "void at::native::unrolled_elementwise_kernel<at::native::FillFunctor<int>, std::array<char*, 1ul>, 4, TrivialOffsetCalculator<0, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithoutCast, at::native::memory::StoreWithoutCast>(int, at::native::FillFunctor<int>, std::array<char*, 1ul>, TrivialOffsetCalculator<0, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithoutCast, at::native::memory::StoreWithoutCast)"
},
{
"duration_ms": 0.0008,
"name": "void at::native::unrolled_elementwise_kernel<at::native::CUDAFunctorOnSelf_add<int>, std::array<char*, 2ul>, 4, TrivialOffsetCalculator<1, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithoutCast, at::native::memory::StoreWithoutCast>(int, at::native::CUDAFunctorOnSelf_add<int>, std::array<char*, 2ul>, TrivialOffsetCalculator<1, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithoutCast, at::native::memory::StoreWithoutCast)"
}
],
"non_kernel_gap_ms": 7.193539999999928,
"selected_execute_annotation": "execute_context_1(8192)_generation_0(0)",
"trace": "runs/frontier-tp2-prefill-serving-v0/fleet-artifacts/tp2-prefill-serving-smoke-20260723-v1-20260723T082454699036Z/artifacts/runs/frontier-tp2-prefill-serving-v0/remote-outputs/tp2-smoke-r1/traces/profile/dp0_pp0_tp1_dcp0_ep1_rank1.1784795312743314215.pt.trace.json.gz"
}
],
"schema": "frontier-tp2-prefill-serving-smoke.v1"
}

View File

@@ -0,0 +1,147 @@
#!/usr/bin/env python3
"""Replay one TP2 trace with structured attention and measured prefill MoE."""
from __future__ import annotations
import argparse
import importlib.util
import json
import subprocess
import sys
from pathlib import Path
ROOT = Path(__file__).resolve().parent
REPO = ROOT.parents[1]
S3_REAL = REPO / "runs/frontier-s3-real-v0"
BASE_REFERENCE = (
REPO
/ "runs/frontier-collective-joint-v0/counterfactual/joint-r2/manifest.json"
)
BASE_COMMIT = "deadc4a321f0baaa534c6ebd17f974123733cdc2"
EXPERIMENT_COMMIT = "1f8900a4ac64e45754b03d0aa7c1dddab65785cf"
ATTENTION_PATCH = (
REPO
/ "runs/frontier-attn-structured-v0/"
"0001-Experiment-with-structured-attention-prefill-predict.patch"
)
VERDICT = ROOT / "results/serving-smoke-verdict.json"
def load_s3_module():
spec = importlib.util.spec_from_file_location(
"s3_prefix_replay", S3_REAL / "run_frontier_prefix_replay.py"
)
module = importlib.util.module_from_spec(spec)
sys.path.insert(0, str(S3_REAL))
spec.loader.exec_module(module)
return module
def git(checkout: Path, *args: str) -> str:
return subprocess.check_output(
["git", "-C", str(checkout), *args], text=True
).strip()
def main() -> None:
parser = argparse.ArgumentParser()
parser.add_argument("--trace", type=Path, required=True)
parser.add_argument("--output-root", type=Path, required=True)
parser.add_argument("--label", required=True)
parser.add_argument("--max-tokens", type=int, required=True)
parser.add_argument("--duration-s", type=float)
parser.add_argument("--cache-root", type=Path, required=True)
parser.add_argument(
"--frontier-checkout",
type=Path,
default=Path("/tmp/frontier-attn-structured-v0"),
)
parser.add_argument(
"--attention-profile",
type=Path,
default=REPO
/ "runs/frontier-prefill-kvgrowth-fix-v0/profiles/"
"profile-v5-kvgrowth/attention.csv",
)
args = parser.parse_args()
args.config = "tp2_mns16"
frontier = args.frontier_checkout.resolve()
profile = args.attention_profile.resolve()
if git(frontier, "rev-parse", "HEAD") != EXPERIMENT_COMMIT:
raise SystemExit(f"unexpected experiment checkout HEAD: {frontier}")
if git(frontier, "rev-parse", "HEAD^") != BASE_COMMIT:
raise SystemExit("experiment commit is not directly based on frozen Frontier")
if git(frontier, "status", "--porcelain"):
raise SystemExit("experiment Frontier checkout must be clean")
if not profile.is_file():
raise SystemExit(f"attention profile missing: {profile}")
smoke = json.loads(VERDICT.read_text())
moe_scale = float(smoke["counterfactual"]["moe_scale"])
reference = json.loads(BASE_REFERENCE.read_text())
reference["frontier_checkout"] = str(frontier)
reference["frontier_commit"] = EXPERIMENT_COMMIT
generated_reference = ROOT / "frontier-reference.json"
generated_reference.write_text(json.dumps(reference, indent=2))
module = load_s3_module()
module.REFERENCE = generated_reference
module.EXPECTED_FRONTIER_COMMIT = EXPERIMENT_COMMIT
original_replace = module.replace_flag
injection_added = False
def replace_and_override(argv: list[str], flag: str, value: str) -> None:
nonlocal injection_added
original_replace(argv, flag, value)
if flag.endswith("trace_file"):
attention_flag = (
"--random_forrest_execution_time_predictor_config_atten_input_file"
)
original_replace(argv, attention_flag, str(profile))
no_cache = (
"--random_forrest_execution_time_predictor_config_no_cache"
)
if no_cache in argv:
argv.remove(no_cache)
if not injection_added:
argv.extend(
(
"--random_forrest_execution_time_predictor_config_"
"moe_grouped_gemm_calibration_scale",
str(moe_scale),
"--random_forrest_execution_time_predictor_config_"
"decode_phase_moe_grouped_gemm_calibration_scale",
"1.0",
)
)
injection_added = True
module.replace_flag = replace_and_override
module.parse_args = lambda: args
module.main()
manifest_path = args.output_root / "manifest.json"
manifest = json.loads(manifest_path.read_text())
manifest.update(
{
"schema": "frontier-tp2-prefill-serving-replay-v1",
"frontier_base_commit": BASE_COMMIT,
"frontier_experiment_commit": EXPERIMENT_COMMIT,
"structured_attention_patch": str(ATTENTION_PATCH.resolve()),
"structured_attention_patch_sha256": module.sha256(ATTENTION_PATCH),
"attention_profile_override": str(profile),
"attention_profile_sha256": module.sha256(profile),
"prefill_moe_grouped_gemm_scale": moe_scale,
"decode_phase_moe_grouped_gemm_scale": 1.0,
"moe_scale_source": str(VERDICT.resolve()),
"moe_scale_source_sha256": module.sha256(VERDICT),
"model_cache_enabled": True,
}
)
manifest_path.write_text(json.dumps(manifest, indent=2))
print(f"TP2 prefill-MoE replay done: {args.output_root}")
if __name__ == "__main__":
main()