From 4f22688bfdf874a919ceee28055a700f2029db1a Mon Sep 17 00:00:00 2001 From: Gahow Wang Date: Thu, 23 Jul 2026 18:08:32 +0800 Subject: [PATCH] Record TP2 prefill serving-path verdict --- .../.gitignore | 1 + .../analyze_replay_verdict.py | 113 ++++ .../analyze_smoke_verdict.py | 80 +++ .../experiment-card.md | 26 +- .../frontier-reference.json | 525 +++++++++++++++ .../results/replay-verdict.json | 631 ++++++++++++++++++ .../results/serving-smoke-verdict.json | 32 + .../results/serving-smoke.json | 463 +++++++++++++ .../run_replay.py | 147 ++++ 9 files changed, 2013 insertions(+), 5 deletions(-) create mode 100644 runs/frontier-tp2-prefill-serving-v0/analyze_replay_verdict.py create mode 100644 runs/frontier-tp2-prefill-serving-v0/analyze_smoke_verdict.py create mode 100644 runs/frontier-tp2-prefill-serving-v0/frontier-reference.json create mode 100644 runs/frontier-tp2-prefill-serving-v0/results/replay-verdict.json create mode 100644 runs/frontier-tp2-prefill-serving-v0/results/serving-smoke-verdict.json create mode 100644 runs/frontier-tp2-prefill-serving-v0/results/serving-smoke.json create mode 100644 runs/frontier-tp2-prefill-serving-v0/run_replay.py diff --git a/runs/frontier-tp2-prefill-serving-v0/.gitignore b/runs/frontier-tp2-prefill-serving-v0/.gitignore index 011358f..5fa37d2 100644 --- a/runs/frontier-tp2-prefill-serving-v0/.gitignore +++ b/runs/frontier-tp2-prefill-serving-v0/.gitignore @@ -1,3 +1,4 @@ fleet-artifacts/ fleet-state/ +replay/ remote-outputs/ diff --git a/runs/frontier-tp2-prefill-serving-v0/analyze_replay_verdict.py b/runs/frontier-tp2-prefill-serving-v0/analyze_replay_verdict.py new file mode 100644 index 0000000..da0a72c --- /dev/null +++ b/runs/frontier-tp2-prefill-serving-v0/analyze_replay_verdict.py @@ -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() diff --git a/runs/frontier-tp2-prefill-serving-v0/analyze_smoke_verdict.py b/runs/frontier-tp2-prefill-serving-v0/analyze_smoke_verdict.py new file mode 100644 index 0000000..e23e70c --- /dev/null +++ b/runs/frontier-tp2-prefill-serving-v0/analyze_smoke_verdict.py @@ -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() diff --git a/runs/frontier-tp2-prefill-serving-v0/experiment-card.md b/runs/frontier-tp2-prefill-serving-v0/experiment-card.md index 843cfa7..d8944a4 100644 --- a/runs/frontier-tp2-prefill-serving-v0/experiment-card.md +++ b/runs/frontier-tp2-prefill-serving-v0/experiment-card.md @@ -1,6 +1,6 @@ # 实验 EXP-TP2-PREFILL-SERVING：TP2 base-prefill residual 是否来自 serving-path MoE -> **状态：** GPU smoke harness 已冻结，待远端执行 +> **状态：** 完成；机制 PASS，global constant injection FAIL > > 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 component ledger 中 MoE=`171.12 ms`，若单独解释 residual 需增至约 `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 preflight（2026-07-23）：** dash1--dash4 均为 8×H20；32 张卡 `memory.used=0 MiB`、`utilization=0%`，无 compute process。dry-run 选择 `dash1:[0,1]`，正式 job 已 pin 到 dash1。 -- **含义：** 待 GPU。 -- **Claim update：** unchanged。 -- **下一步：** probe fleet；运行 TP2 serving-path prefill smoke。 +- **含义：** TP2 base residual 的主要机制确实是 serving-path MoE，而不是 + attention 或 real-anchor 噪声；但单个 q8k 点不能外推成全 prefill-domain + 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 scale；campaign 继续实验 3。 diff --git a/runs/frontier-tp2-prefill-serving-v0/frontier-reference.json b/runs/frontier-tp2-prefill-serving-v0/frontier-reference.json new file mode 100644 index 0000000..bf87302 --- /dev/null +++ b/runs/frontier-tp2-prefill-serving-v0/frontier-reference.json @@ -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" + } + } +} \ No newline at end of file diff --git a/runs/frontier-tp2-prefill-serving-v0/results/replay-verdict.json b/runs/frontier-tp2-prefill-serving-v0/results/replay-verdict.json new file mode 100644 index 0000000..54011a6 --- /dev/null +++ b/runs/frontier-tp2-prefill-serving-v0/results/replay-verdict.json @@ -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" +} diff --git a/runs/frontier-tp2-prefill-serving-v0/results/serving-smoke-verdict.json b/runs/frontier-tp2-prefill-serving-v0/results/serving-smoke-verdict.json new file mode 100644 index 0000000..2b7f3a3 --- /dev/null +++ b/runs/frontier-tp2-prefill-serving-v0/results/serving-smoke-verdict.json @@ -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" +} diff --git a/runs/frontier-tp2-prefill-serving-v0/results/serving-smoke.json b/runs/frontier-tp2-prefill-serving-v0/results/serving-smoke.json new file mode 100644 index 0000000..ccf2f25 --- /dev/null +++ b/runs/frontier-tp2-prefill-serving-v0/results/serving-smoke.json @@ -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::Params)" + }, + { + "duration_ms": 97.15557600000001, + "name": "void cutlass::device_kernel, cute::C<1>, cute::C<1> >, cute::tuple, 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::C<128>, cute::C<128> >, cute::tuple, 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, cute::C<1>, cute::C<1> >, cute::tuple, 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::C<128>, cute::C<128> >, cute::tuple, 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, std::array >(int, at::native::FillFunctor, std::array)" + }, + { + "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, 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, 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::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef, at::native::index_kernel_impl >(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef)::{lambda(char*, char const*, long)#1} const&, bool)::{lambda(int)#1}>(long, at::native::gpu_index_kernel >(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef, at::native::index_kernel_impl >(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef)::{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::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef, at::native::index_kernel_impl >(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef)::{lambda(char*, char const*, long)#1} const&, bool)::{lambda(int)#1}>(long, at::native::gpu_index_kernel >(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef, at::native::index_kernel_impl >(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef)::{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, std::array >(int, at::native::CUDAFunctor_add, std::array)" + }, + { + "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, std::array, 4, TrivialOffsetCalculator<0, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithoutCast, at::native::memory::StoreWithoutCast>(int, at::native::FillFunctor, std::array, 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, std::array >(int, at::native::CUDAFunctor_add, std::array)" + }, + { + "duration_ms": 0.000767, + "name": "void at::native::unrolled_elementwise_kernel, std::array, 4, TrivialOffsetCalculator<1, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithoutCast, at::native::memory::StoreWithoutCast>(int, at::native::CUDAFunctorOnSelf_add, std::array, 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::Params)" + }, + { + "duration_ms": 97.15557600000001, + "name": "void cutlass::device_kernel, cute::C<1>, cute::C<1> >, cute::tuple, 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::C<128>, cute::C<128> >, cute::tuple, 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, cute::C<1>, cute::C<1> >, cute::tuple, 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::C<128>, cute::C<128> >, cute::tuple, 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, std::array >(int, at::native::FillFunctor, std::array)" + }, + { + "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, 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, 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::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef, at::native::index_kernel_impl >(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef)::{lambda(char*, char const*, long)#1} const&, bool)::{lambda(int)#1}>(long, at::native::gpu_index_kernel >(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef, at::native::index_kernel_impl >(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef)::{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::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef, at::native::index_kernel_impl >(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef)::{lambda(char*, char const*, long)#1} const&, bool)::{lambda(int)#1}>(long, at::native::gpu_index_kernel >(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef, at::native::index_kernel_impl >(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef)::{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, std::array >(int, at::native::CUDAFunctor_add, std::array)" + }, + { + "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, std::array, 4, TrivialOffsetCalculator<0, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithoutCast, at::native::memory::StoreWithoutCast>(int, at::native::FillFunctor, std::array, 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, std::array >(int, at::native::CUDAFunctor_add, std::array)" + }, + { + "duration_ms": 0.000767, + "name": "void at::native::unrolled_elementwise_kernel, std::array, 4, TrivialOffsetCalculator<1, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithoutCast, at::native::memory::StoreWithoutCast>(int, at::native::CUDAFunctorOnSelf_add, std::array, 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::Params)" + }, + { + "duration_ms": 97.24454899999998, + "name": "void cutlass::device_kernel, cute::C<1>, cute::C<1> >, cute::tuple, 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::C<128>, cute::C<128> >, cute::tuple, 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, cute::C<1>, cute::C<1> >, cute::tuple, 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::C<128>, cute::C<128> >, cute::tuple, 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, std::array >(int, at::native::FillFunctor, std::array)" + }, + { + "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, 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, 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::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef, at::native::index_kernel_impl >(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef)::{lambda(char*, char const*, long)#1} const&, bool)::{lambda(int)#1}>(long, at::native::gpu_index_kernel >(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef, at::native::index_kernel_impl >(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef)::{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::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef, at::native::index_kernel_impl >(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef)::{lambda(char*, char const*, long)#1} const&, bool)::{lambda(int)#1}>(long, at::native::gpu_index_kernel >(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef)::{lambda(char*, char const*, long)#1}>(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef, at::native::index_kernel_impl >(at::TensorIteratorBase&, c10::ArrayRef, c10::ArrayRef)::{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, std::array >(int, at::native::CUDAFunctor_add, std::array)" + }, + { + "duration_ms": 0.001184, + "name": "void at::native::vectorized_elementwise_kernel<4, at::native::CUDAFunctor_add, std::array >(int, at::native::CUDAFunctor_add, std::array)" + }, + { + "duration_ms": 0.001184, + "name": "void at::native::unrolled_elementwise_kernel, std::array, 4, TrivialOffsetCalculator<0, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithoutCast, at::native::memory::StoreWithoutCast>(int, at::native::FillFunctor, std::array, 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, std::array, 4, TrivialOffsetCalculator<1, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithoutCast, at::native::memory::StoreWithoutCast>(int, at::native::CUDAFunctorOnSelf_add, std::array, 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" +} diff --git a/runs/frontier-tp2-prefill-serving-v0/run_replay.py b/runs/frontier-tp2-prefill-serving-v0/run_replay.py new file mode 100644 index 0000000..dee3c04 --- /dev/null +++ b/runs/frontier-tp2-prefill-serving-v0/run_replay.py @@ -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()