From 5cb9a89dc341c5053fe4c9e483f9bbce1280c228 Mon Sep 17 00:00:00 2001 From: flyingdown Date: Wed, 5 Aug 2026 05:47:33 +0800 Subject: [PATCH 01/15] feat(sglang-trace-analyze): scaffold new task plugin for SGLang torch profiler analysis Adds a new task type that profiles models across multiple batch sizes using sglang's bench_one_batch_server with torch profiler, then analyzes the traces for kernel hotspots, TFLOPS/MFU, operator-to-model-structure mapping, fuse opportunities, and LLM-powered optimization hints. 5-phase linear pipeline: MAPPING -> BENCHMARK -> ANALYZE -> HINTS -> SUMMARIZE. Design doc: docs/sglang_trace_analyze-design.md (grilled by architecture review). Co-Authored-By: deepseek-v4-pro[1m] --- .../tasks/sglang_trace_analyze/__init__.py | 8 + .../tasks/sglang_trace_analyze/form.yaml | 67 +++ .../orchestrator/__init__.py | 6 + .../sglang_trace_analyze/orchestrator/cli.py | 46 ++ .../orchestrator/flops_calculator.py | 175 ++++++ .../orchestrator/fuse_matcher.py | 136 +++++ .../orchestrator/gpu_specs.py | 63 ++ .../orchestrator/iteration_record.py | 210 +++++++ .../orchestrator/orchestrator.py | 75 +++ .../orchestrator/overlap_detector.py | 115 ++++ .../orchestrator/phases.py | 55 ++ .../orchestrator/pipeline.py | 564 ++++++++++++++++++ .../orchestrator/plugin.py | 10 + .../orchestrator/prompts.py | 117 ++++ .../orchestrator/run_benchmark.py | 152 +++++ .../orchestrator/structure_mapper.py | 158 +++++ .../orchestrator/trace_parser.py | 140 +++++ .../sglang_trace_analyze/server/__init__.py | 1 + .../server/_state_readers.py | 51 ++ .../sglang_trace_analyze/server/plugin.py | 31 + .../sglang_trace_analyze/server/routes.py | 68 +++ .../sglang_trace_analyze/tests/__init__.py | 0 .../tests/test_flops_calculator.py | 78 +++ .../tests/test_fuse_matcher.py | 37 ++ .../tests/test_gpu_specs.py | 17 + .../sglang_trace_analyze/tests/test_plugin.py | 18 + .../tests/test_server_readers.py | 75 +++ .../tests/test_structure_mapper.py | 63 ++ .../tests/test_trace_parser.py | 66 ++ 29 files changed, 2602 insertions(+) create mode 100644 metainfer/tasks/sglang_trace_analyze/__init__.py create mode 100644 metainfer/tasks/sglang_trace_analyze/form.yaml create mode 100644 metainfer/tasks/sglang_trace_analyze/orchestrator/__init__.py create mode 100644 metainfer/tasks/sglang_trace_analyze/orchestrator/cli.py create mode 100644 metainfer/tasks/sglang_trace_analyze/orchestrator/flops_calculator.py create mode 100644 metainfer/tasks/sglang_trace_analyze/orchestrator/fuse_matcher.py create mode 100644 metainfer/tasks/sglang_trace_analyze/orchestrator/gpu_specs.py create mode 100644 metainfer/tasks/sglang_trace_analyze/orchestrator/iteration_record.py create mode 100644 metainfer/tasks/sglang_trace_analyze/orchestrator/orchestrator.py create mode 100644 metainfer/tasks/sglang_trace_analyze/orchestrator/overlap_detector.py create mode 100644 metainfer/tasks/sglang_trace_analyze/orchestrator/phases.py create mode 100644 metainfer/tasks/sglang_trace_analyze/orchestrator/pipeline.py create mode 100644 metainfer/tasks/sglang_trace_analyze/orchestrator/plugin.py create mode 100644 metainfer/tasks/sglang_trace_analyze/orchestrator/prompts.py create mode 100644 metainfer/tasks/sglang_trace_analyze/orchestrator/run_benchmark.py create mode 100644 metainfer/tasks/sglang_trace_analyze/orchestrator/structure_mapper.py create mode 100644 metainfer/tasks/sglang_trace_analyze/orchestrator/trace_parser.py create mode 100644 metainfer/tasks/sglang_trace_analyze/server/__init__.py create mode 100644 metainfer/tasks/sglang_trace_analyze/server/_state_readers.py create mode 100644 metainfer/tasks/sglang_trace_analyze/server/plugin.py create mode 100644 metainfer/tasks/sglang_trace_analyze/server/routes.py create mode 100644 metainfer/tasks/sglang_trace_analyze/tests/__init__.py create mode 100644 metainfer/tasks/sglang_trace_analyze/tests/test_flops_calculator.py create mode 100644 metainfer/tasks/sglang_trace_analyze/tests/test_fuse_matcher.py create mode 100644 metainfer/tasks/sglang_trace_analyze/tests/test_gpu_specs.py create mode 100644 metainfer/tasks/sglang_trace_analyze/tests/test_plugin.py create mode 100644 metainfer/tasks/sglang_trace_analyze/tests/test_server_readers.py create mode 100644 metainfer/tasks/sglang_trace_analyze/tests/test_structure_mapper.py create mode 100644 metainfer/tasks/sglang_trace_analyze/tests/test_trace_parser.py diff --git a/metainfer/tasks/sglang_trace_analyze/__init__.py b/metainfer/tasks/sglang_trace_analyze/__init__.py new file mode 100644 index 00000000..afb0e3af --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/__init__.py @@ -0,0 +1,8 @@ +"""sglang_trace_analyze — auto-generate torch profiler traces via SGLang, +analyze them (operator-to-structure mapping, kernel hotspots, TFLOPS / MFU, +overlap opportunities, fuse suggestions), and surface results + LLM hints +in the MetaInfer WebUI. +""" + +from .orchestrator import plugin as _task_plugin # noqa: F401 +from .server import plugin as _web_plugin # noqa: F401 diff --git a/metainfer/tasks/sglang_trace_analyze/form.yaml b/metainfer/tasks/sglang_trace_analyze/form.yaml new file mode 100644 index 00000000..fdd2dc1b --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/form.yaml @@ -0,0 +1,67 @@ +- key: model_path + header: Model Path + question: "HuggingFace repo id or local path to the model weights." + required: true + form: text + +- key: version + header: Version + question: "Short tag for this run — used in trace directory naming and result labels." + required: true + form: text + +- key: batch_sizes + header: Batch Sizes + question: "Comma-separated list of decode batch sizes to profile, e.g. 1,4,8,16." + required: true + form: text + +- key: mapping_batch_size + header: Mapping BS + question: "Batch size for the mapping run (CUDA Graph disabled). One value is enough — kernel-to-layer mapping is independent of batch size." + required: true + default: "8" + form: number + +- key: input_len + header: Input Len + question: "Synthetic input sequence length." + required: true + default: "512" + form: number + +- key: output_len + header: Output Len + question: "Synthetic output sequence length." + required: true + default: "2000" + form: number + +- key: tp_size + header: TP Size + question: "Tensor-parallelism degree." + required: true + default: "1" + form: number + +- key: pp_size + header: PP Size + question: "Pipeline-parallelism degree." + required: true + default: "1" + form: number + +- key: gpu_model + header: GPU Model + question: "GPU model — used to look up theoretical peak TFLOPS and memory bandwidth." + required: true + form: select + options: + - label: "K100" + description: "FP32 49TF, TF32 98TF, BF16/FP16 192TF, INT8 392TOPS, BW 700GB/s" + - label: "A100_80G" + description: "FP32 19.5TF, TF32 156TF, BF16/FP16 312TF, INT8 624TOPS, BW 2039GB/s" + - label: "H100" + description: "FP32 67TF, TF32 989TF, BF16/FP16 989TF, INT8 1979TOPS, BW 3350GB/s" + - label: "B200" + description: "FP32 90TF, TF32 2250TF, BF16/FP16 2250TF, INT8 4500TOPS, BW 8000GB/s" diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/__init__.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/__init__.py new file mode 100644 index 00000000..43d16d09 --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/__init__.py @@ -0,0 +1,6 @@ +"""Orchestrator (worker subprocess) for sglang_trace_analyze.""" + +from metainfer.orchestrator.tasks import register +from .plugin import PLUGIN + +register(PLUGIN) diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/cli.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/cli.py new file mode 100644 index 00000000..d4179885 --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/cli.py @@ -0,0 +1,46 @@ +"""CLI entry point for the sglang_trace_analyze orchestrator subprocess. + +The launcher spawns:: + + python -m run --state-dir … --workspace-dir … + +Contract required by the framework (§6d): ``run`` subcommand + ``--state-dir`` +and ``--workspace-dir`` flags. +""" + +from __future__ import annotations + +import argparse +import sys +from pathlib import Path + + +def main(argv: list[str] | None = None) -> int: + parser = argparse.ArgumentParser(prog="metainfer-orchestrator") + sub = parser.add_subparsers(dest="command") + + run_p = sub.add_parser("run") + run_p.add_argument("requirements", type=Path, + help="Path to requirements.json") + run_p.add_argument("--state-dir", type=Path, required=True) + run_p.add_argument("--workspace-dir", type=Path, required=True) + # Task-specific flags + run_p.add_argument("--iter-limit", type=int, default=None, + help="Override max iterations (default: derive from batch count)") + + args = parser.parse_args(argv) + if args.command != "run": + parser.print_help() + return 1 + + from .orchestrator import run_with_requirements + return run_with_requirements( + requirements_path=args.requirements, + state_dir=args.state_dir, + workspace_dir=args.workspace_dir, + iter_limit=args.iter_limit, + ) + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/flops_calculator.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/flops_calculator.py new file mode 100644 index 00000000..e4192a41 --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/flops_calculator.py @@ -0,0 +1,175 @@ +"""Compute TFLOPS, bandwidth, and MFU for aggregated kernel entries. + +Uses: +- ``gpu_specs.py`` for theoretical peak values +- kernel ``input_dims`` (from MAPPING trace) or shape rules (for CUDA Graph + formal traces) to derive actual FLOP counts per invocation +- kernel ``total_dur_us`` to compute actual TFLOPS/bandwidth +""" + +from __future__ import annotations + +from typing import Any, Dict, List, Optional + +from .gpu_specs import GpuSpec + + +def calculate_mfu( + kernels: List[Dict[str, Any]], + gpu_spec: GpuSpec, + *, + batch_size: int, + dtype: str = "bf16", +) -> List[Dict[str, Any]]: + """Augment each kernel entry with TFLOPS, bandwidth, MFU, and bound classification. + + Args: + kernels: Aggregated kernel list. Each entry must have ``total_dur_us`` + and ``count``. Entries from a non-CUDA Graph trace may also have + ``input_dims``, which are used for FLOP/byte estimation where available. + gpu_spec: GPU theoretical peak specification. + batch_size: Decode batch size used for this trace. + dtype: Compute dtype — determines which TFLOPS peak to use. + One of ``fp32``, ``tf32``, ``bf16``, ``fp16``, ``int8``. + + Returns: + The same kernel list with added fields: ``tflops_actual``, + ``bandwidth_gb_s``, ``mfu``, ``bound``, ``flops_per_invocation``. + """ + theoretical_tflops = _theoretical_peak(gpu_spec, dtype) + theoretical_bw = gpu_spec.bandwidth_gb_s + + for k in kernels: + dur_s = k["total_dur_us"] / 1e6 + count = k.get("count", 1) + dur_per_invocation_s = dur_s / count if count else dur_s + dims = k.get("input_dims", []) + op_type = k.get("op_type", "Other") + + flops = _estimate_flops(op_type, dims, batch_size) + bytes_moved = _estimate_bytes(op_type, dims, batch_size) + + tflops_actual = (flops / dur_s / 1e12) if dur_s > 0 else 0 + bandwidth_gb_s = (bytes_moved / dur_s / 1e9) if dur_s > 0 else 0 + mfu = (tflops_actual / theoretical_tflops * 100) if theoretical_tflops > 0 else 0 + + # Compute-bound vs memory-bound heuristic + ops_per_byte = flops / bytes_moved if bytes_moved > 0 else float("inf") + # "Roofline" crossover point = peak_flops / peak_bw ops/byte + if theoretical_bw > 0: + crossover = theoretical_tflops * 1e12 / (theoretical_bw * 1e9) + else: + crossover = float("inf") + bound = "compute" if ops_per_byte > crossover else "memory" + + k["tflops_actual"] = round(tflops_actual, 3) + k["tflops_theoretical"] = theoretical_tflops + k["bandwidth_gb_s"] = round(bandwidth_gb_s, 1) + k["bandwidth_theoretical"] = theoretical_bw + k["mfu"] = round(mfu, 1) + k["bound"] = bound + k["flops_per_invocation"] = int(flops) + + return kernels + + +def _theoretical_peak(spec: GpuSpec, dtype: str) -> float: + """Return theoretical peak TFLOPS for the given dtype.""" + return { + "fp32": spec.fp32_tflops, + "tf32": spec.tf32_tflops, + "bf16": spec.bf16_tflops, + "fp16": spec.fp16_tflops, + "int8": spec.int8_tops, # TOPS → TFLOPS approximate + }.get(dtype, spec.bf16_tflops) + + +def _estimate_flops( + op_type: str, + dims: List[Any], + batch_size: int, +) -> float: + """Estimate FLOPs for one kernel invocation. + + For GEMM: 2*M*N*K (or 2*B*M*N*K for batched). + For Attention: approximately 4*B*seq_len*head_dim*num_heads^2. + For ElementWise: 2*num_elements. + + Returns 0 if dims are unavailable (CUDA Graph trace). + """ + if not dims: + return 0 + + # Use the first observed dim list + d = dims[0] + + if op_type == "GEMM": + if isinstance(d, list) and len(d) >= 2: + if len(d) == 3: + M, K, N = int(d[0]), int(d[1]), int(d[2]) + return 2 * M * N * K + B, M, N, K = _unpack_4d(d, batch_size) + return 2 * B * M * N * K + + elif op_type == "Attention": + if isinstance(d, list) and len(d) >= 3: + seq_len = int(d[0]) + num_heads = int(d[1]) + head_dim = int(d[2]) + return 4 * seq_len * head_dim * num_heads * num_heads * batch_size + + elif op_type == "MoE": + if isinstance(d, list) and len(d) >= 3: + M, K, N = int(d[0]), int(d[1]), int(d[2]) + return 2 * M * N * K + + return 0 + + +def _estimate_bytes( + op_type: str, + dims: List[Any], + batch_size: int, +) -> float: + """Estimate bytes moved (reads + writes) for one kernel invocation. + + Simple heuristic: for GEMM, input_bytes ≈ (M*K + K*N) * dtype_size, + output_bytes ≈ M*N * dtype_size. For elementwise, ≈ 3 * num_elements. + + Returns 0 if dims are unavailable. + """ + if not dims: + return 0 + + d = dims[0] + dtype_size = 2 # bf16/fp16 default + + if op_type == "GEMM": + if isinstance(d, list): + if len(d) == 3: + M, K, N = int(d[0]), int(d[1]), int(d[2]) + return (M * K + K * N + M * N) * dtype_size + B, M, N, K = _unpack_4d(d, batch_size) + return B * (M * K + K * N + M * N) * dtype_size + + elif op_type == "Attention": + if isinstance(d, list) and len(d) >= 3: + seq_len = int(d[0]) + num_heads = int(d[1]) + head_dim = int(d[2]) + # Q, K, V reads + output write (approximate) + return batch_size * seq_len * num_heads * head_dim * 4 * dtype_size + + return 0 + + +def _unpack_4d( + dims: list, + batch_size: int, +) -> tuple: + """Unpack a 4-element dim list into (B, M, N, K), defaulting B to batch_size.""" + if len(dims) >= 4: + return int(dims[0]), int(dims[1]), int(dims[2]), int(dims[3]) + if len(dims) == 3: + return batch_size, int(dims[0]), int(dims[1]), int(dims[2]) + return batch_size, int(dims[0]), 1, 1 diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/fuse_matcher.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/fuse_matcher.py new file mode 100644 index 00000000..58b21ffe --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/fuse_matcher.py @@ -0,0 +1,136 @@ +"""Rule-based fuse pattern matcher. + +Scans the kernel table (ordered by GPU time or timeline order) for known +sequences that indicate a missing fusion opportunity, and reports each +match with a description and estimated saving. + +The catalog is hard-coded — each pattern has a name, the kernel names +that must appear consecutively (or within a short window), and a +suggestion for what the fused replacement would be. +""" + +from __future__ import annotations + +from typing import Any, Dict, List + +# ------------------------------------------------------------------ # +# Fuse pattern catalog +# ------------------------------------------------------------------ # + +FUSE_PATTERNS: List[Dict[str, Any]] = [ + { + "pattern": "rms_norm + gemm", + "kernels": ["rms_norm", "gemm"], + "match_mode": "consecutive", + "suggestion": "Replace separate rms_norm + gemm with fused_rms_norm_gemm (e.g. triton kernel or sglang fused op).", + "estimated_saving_us": 180, + "confidence": "high", + }, + { + "pattern": "silu + mul + gemm", + "kernels": ["silu", "mul", "gemm"], + "match_mode": "consecutive", + "suggestion": "Fuse into silu_and_mul + gemm, or a single fused MoE activation+gemm kernel.", + "estimated_saving_us": 250, + "confidence": "high", + }, + { + "pattern": "add + rms_norm", + "kernels": ["add", "rms_norm"], + "match_mode": "consecutive", + "suggestion": "Fuse residual add + rms_norm into a single kernel to avoid a separate memory round-trip.", + "estimated_saving_us": 120, + "confidence": "medium", + }, + { + "pattern": "quant + gemm", + "kernels": ["quant", "gemm"], + "match_mode": "consecutive", + "suggestion": "Integrate FP8 quantization into the GEMM launch to eliminate a precursor kernel.", + "estimated_saving_us": 200, + "confidence": "medium", + }, + { + "pattern": "nccl_allreduce + gemm (no overlap)", + "kernels": ["ncclAllReduce", "gemm"], + "match_mode": "consecutive", + "suggestion": ( + "AllReduce and gemm are serialized. Try overlapping: issue AllReduce " + "on a separate CUDA stream, or restructure to compute on one output " + "shard while communicating another." + ), + "estimated_saving_us": 300, + "confidence": "medium", + }, +] + + +def match_fuse_patterns( + kernels: List[Dict[str, Any]], +) -> List[Dict[str, Any]]: + """Scan a kernel list for known fuse patterns. + + Args: + kernels: List of kernel entries. Must contain ``kernel_name`` and + preferably be in timeline order. If only duration-ordered, set + ``match_mode`` to ``"unordered"`` for pattern matching. + + Returns: + List of matched patterns, each with ``pattern``, ``kernels``, + ``suggestion``, ``estimated_saving_us``, ``confidence``. + """ + kernel_names = [k.get("kernel_name", "") for k in kernels] + matches = [] + + for pat in FUSE_PATTERNS: + found = _match_consecutive(kernel_names, pat["kernels"]) + if found: + matches.append({ + "pattern": pat["pattern"], + "kernels": found, + "suggestion": pat["suggestion"], + "estimated_saving_us": pat["estimated_saving_us"], + "confidence": pat["confidence"], + }) + + return matches + + +def build_fuse_report( + kernels: List[Dict[str, Any]], + batch_size: int, + stage: str, +) -> Dict[str, Any]: + """Produce the full fuse.json payload.""" + matches = match_fuse_patterns(kernels) + return { + "batch_size": batch_size, + "stage": stage, + "matches": matches, + } + + +def _match_consecutive( + names: List[str], + pattern_kernels: List[str], +) -> List[str]: + """Check if ``pattern_kernels`` appear consecutively (in order) within + ``names``. + + Returns the matched kernel names if found, empty list otherwise. + """ + if len(pattern_kernels) > len(names): + return [] + + patterns_lower = [p.lower() for p in pattern_kernels] + names_lower = [n.lower() for n in names] + + for i in range(len(names_lower) - len(patterns_lower) + 1): + match = True + for j, pat in enumerate(patterns_lower): + if pat not in names_lower[i + j]: + match = False + break + if match: + return names[i: i + len(patterns_lower)] + return [] diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/gpu_specs.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/gpu_specs.py new file mode 100644 index 00000000..d88abef7 --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/gpu_specs.py @@ -0,0 +1,63 @@ +"""GPU theoretical-peak lookup table. + +Used by ``flops_calculator.py`` to compute MFU: + MFU = actual_TFLOPS / theoretical_peak_TFLOPS. +""" + +from __future__ import annotations + +from dataclasses import dataclass +from typing import Dict + + +@dataclass(frozen=True) +class GpuSpec: + """Theoretical peak numbers for one GPU model.""" + + label: str + fp32_tflops: float + tf32_tflops: float + bf16_tflops: float + fp16_tflops: float + int8_tops: float + bandwidth_gb_s: float + + +GPU_SPECS: Dict[str, GpuSpec] = { + "K100": GpuSpec( + label="K100", + fp32_tflops=49, + tf32_tflops=98, + bf16_tflops=192, + fp16_tflops=192, + int8_tops=392, + bandwidth_gb_s=700, + ), + "A100_80G": GpuSpec( + label="A100_80G", + fp32_tflops=19.5, + tf32_tflops=156, + bf16_tflops=312, + fp16_tflops=312, + int8_tops=624, + bandwidth_gb_s=2039, + ), + "H100": GpuSpec( + label="H100", + fp32_tflops=67, + tf32_tflops=989, + bf16_tflops=989, + fp16_tflops=989, + int8_tops=1979, + bandwidth_gb_s=3350, + ), + "B200": GpuSpec( + label="B200", + fp32_tflops=90, + tf32_tflops=2250, + bf16_tflops=2250, + fp16_tflops=2250, + int8_tops=4500, + bandwidth_gb_s=8000, + ), +} diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/iteration_record.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/iteration_record.py new file mode 100644 index 00000000..c014831a --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/iteration_record.py @@ -0,0 +1,210 @@ +"""Phase-specific iteration records for sglang_trace_analyze. + +Each phase gets its own dataclass so the schema stays clean — no +``None``-filled optional fields bleeding across phases. +""" + +from __future__ import annotations + +import time +from dataclasses import dataclass, field, fields +from typing import Any, Dict + + +def _base_dict(rec, **overrides) -> Dict[str, Any]: + """Serialize *any* iteration record to a dict the WebUI can read. + + Keys: phase (str), status, started_at, ended_at, plus phase-specific + fields from the dataclass. + """ + out: Dict[str, Any] = { + "phase": getattr(rec, "phase", ""), + "status": rec.status, + "started_at": rec.started_at, + "ended_at": rec.ended_at, + } + for f in fields(rec): + if f.name in ("phase", "status", "started_at", "ended_at"): + continue + val = getattr(rec, f.name) + if val is not None: + out[f.name] = val + out.update(overrides) + return out + + +# ------------------------------------------------------------------ # +# MAPPING phase +# ------------------------------------------------------------------ # + +@dataclass +class MappingRecord: + phase: str = "mapping" + status: str = "running" + started_at: float = 0.0 + ended_at: float = 0.0 + batch_size: int | None = None + trace_dir: str | None = None + duration_s: float | None = None + kernel_count: int | None = None + confidence_issues: int = 0 # entries with low confidence after LLM check + error: str | None = None + + def start(self): + self.started_at = time.time() + self.status = "running" + + def done(self, **kw): + self.status = "success" + self.ended_at = time.time() + for k, v in kw.items(): + setattr(self, k, v) + + def fail(self, error: str): + self.status = "failed" + self.ended_at = time.time() + self.error = error + + def to_dict(self) -> Dict[str, Any]: + return _base_dict(self) + + +# ------------------------------------------------------------------ # +# BENCHMARK phase +# ------------------------------------------------------------------ # + +@dataclass +class BenchmarkRecord: + phase: str = "benchmark" + status: str = "running" + started_at: float = 0.0 + ended_at: float = 0.0 + batch_size: int | None = None + trace_dir: str | None = None + duration_s: float | None = None + throughput: float | None = None + latency_p50: float | None = None + error: str | None = None + + def start(self): + self.started_at = time.time() + self.status = "running" + + def done(self, **kw): + self.status = "success" + self.ended_at = time.time() + for k, v in kw.items(): + setattr(self, k, v) + + def fail(self, error: str): + self.status = "failed" + self.ended_at = time.time() + self.error = error + + def to_dict(self) -> Dict[str, Any]: + return _base_dict(self) + + +# ------------------------------------------------------------------ # +# ANALYZE phase +# ------------------------------------------------------------------ # + +@dataclass +class AnalyzeRecord: + phase: str = "analyze" + status: str = "running" + started_at: float = 0.0 + ended_at: float = 0.0 + batch_size: int | None = None + stage: str | None = None # "prefill" | "decode" + kernel_count: int | None = None + top_kernel: str | None = None + top_kernel_pct: float | None = None + mfu_avg: float | None = None + fuse_hits: int = 0 + error: str | None = None + + def start(self): + self.started_at = time.time() + self.status = "running" + + def done(self, **kw): + self.status = "success" + self.ended_at = time.time() + for k, v in kw.items(): + setattr(self, k, v) + + def fail(self, error: str): + self.status = "failed" + self.ended_at = time.time() + self.error = error + + def to_dict(self) -> Dict[str, Any]: + return _base_dict(self) + + +# ------------------------------------------------------------------ # +# HINTS phase +# ------------------------------------------------------------------ # + +@dataclass +class HintsRecord: + phase: str = "hints" + status: str = "running" + started_at: float = 0.0 + ended_at: float = 0.0 + model_used: str | None = None + batch_count: int = 0 + error: str | None = None + + def start(self): + self.started_at = time.time() + self.status = "running" + + def done(self, **kw): + self.status = "success" + self.ended_at = time.time() + for k, v in kw.items(): + setattr(self, k, v) + + def fail(self, error: str): + self.status = "failed" + self.ended_at = time.time() + self.error = error + + def to_dict(self) -> Dict[str, Any]: + return _base_dict(self) + + +# ------------------------------------------------------------------ # +# SUMMARIZE phase +# ------------------------------------------------------------------ # + +@dataclass +class SummarizeRecord: + phase: str = "summarize" + status: str = "running" + started_at: float = 0.0 + ended_at: float = 0.0 + batch_count: int = 0 + best_batch: int | None = None + best_mfu: float | None = None + error: str | None = None + + def start(self): + self.started_at = time.time() + self.status = "running" + + def done(self, **kw): + self.status = "success" + self.ended_at = time.time() + for k, v in kw.items(): + setattr(self, k, v) + + def fail(self, error: str): + self.status = "failed" + self.ended_at = time.time() + self.error = error + + def to_dict(self) -> Dict[str, Any]: + return _base_dict(self) diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/orchestrator.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/orchestrator.py new file mode 100644 index 00000000..28f70a9e --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/orchestrator.py @@ -0,0 +1,75 @@ +"""Bootstrap + entry point for the sglang_trace_analyze orchestrator. + +Spawns as a child of the WebUI server per task. Reads requirements, sets +up state_dir / workspace_dir, and runs the linear phase pipeline: + + MAPPING -> BENCHMARK -> ANALYZE -> HINTS -> SUMMARIZE -> done + +Unlike gen_infer_framework, this task has no complex transition table — +just five sequential phases with per-(bs, stage) iterations inside +ANALYZE. +""" + +from __future__ import annotations + +import json +from pathlib import Path +from typing import Any, Dict, Optional + +from .pipeline import Pipeline +from .iteration_record import AnalyzeRecord +from metainfer.orchestrator.state import StateStore + + +def run_with_requirements( + requirements_path: Path, + *, + state_dir: Optional[Path] = None, + workspace_dir: Optional[Path] = None, + iter_limit: Optional[int] = None, +) -> int: + """Per-task orchestrator entry point. + + Reads ``requirements.json``, runs the five-phase pipeline to + completion, and exits. + """ + if not requirements_path.exists(): + raise FileNotFoundError(f"requirements file not found: {requirements_path}") + + req: Dict[str, Any] = json.loads( + requirements_path.read_text(encoding="utf-8") + ) + task_id = req.get("task_id", "task") + + # Resolve state_dir + workspace_dir + if state_dir is None or workspace_dir is None: + from metainfer.server import paths as _web_paths + if state_dir is None: + state_dir = _web_paths.task_dir(task_id) + if workspace_dir is None: + workspace_dir = _web_paths.workspace_dir(task_id) + + state_dir.mkdir(parents=True, exist_ok=True) + workspace_dir.mkdir(parents=True, exist_ok=True) + + # Copy requirements into state_dir for self-containment + target_req = state_dir / "requirements.json" + if requirements_path.resolve() != target_req.resolve(): + target_req.write_text( + requirements_path.read_text(encoding="utf-8"), encoding="utf-8" + ) + + store = StateStore(state_dir) + pipe = Pipeline( + req=req, + store=store, + state_dir=state_dir, + workspace_dir=workspace_dir, + ) + + print(f"[metainfer:sglang_trace_analyze] task_id = {task_id}") + print(f"[metainfer:sglang_trace_analyze] state_dir = {state_dir}") + print(f"[metainfer:sglang_trace_analyze] workspace_dir = {workspace_dir}") + + pipe.run() + return 0 diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/overlap_detector.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/overlap_detector.py new file mode 100644 index 00000000..e652d66e --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/overlap_detector.py @@ -0,0 +1,115 @@ +"""Detect communication-computation overlap gaps in a torch profiler trace. + +Scans GPU kernel timeline for gaps between consecutive events where the +GPU is idle. On K100, this is lower priority — the detector is kept +simple. +""" + +from __future__ import annotations + +from typing import Any, Dict, List + + +def detect_gaps( + trace_data: Dict[str, Any], + *, + gap_threshold_us: float = 10.0, +) -> List[Dict[str, Any]]: + """Find GPU-idle gaps in the kernel timeline. + + Args: + trace_data: Parsed Chrome trace JSON. + gap_threshold_us: Minimum gap duration (us) to report. + + Returns: + List of gap dicts with ``gap_id``, ``description``, ``gap_us``, + ``affected_kernels``, ``severity``. + """ + trace_events = trace_data.get("traceEvents", []) + if isinstance(trace_data, list): + trace_events = trace_data + + # Collect GPU kernel events with their timestamps + events = [] + for evt in trace_events: + cat = evt.get("cat", "") + dur = evt.get("dur", 0) + ts = evt.get("ts", 0) + if cat == "kernel" and dur > 0: + events.append({ + "name": evt.get("name", ""), + "ts": ts, + "end": ts + dur, + }) + + events.sort(key=lambda e: e["ts"]) + + gaps = [] + gap_id = 0 + for i in range(1, len(events)): + prev_end = events[i - 1]["end"] + curr_start = events[i]["ts"] + gap = curr_start - prev_end + if gap > gap_threshold_us: + gap_id += 1 + severity = "low" + if gap > 100: + severity = "high" + elif gap > 50: + severity = "medium" + + gaps.append({ + "gap_id": gap_id, + "description": ( + f"{events[i - 1]['name']} → {events[i]['name']}: " + f"{gap:.1f}us idle" + ), + "gap_us": round(gap, 1), + "cumulative_gap_us": 0, # filled in by caller + "pct_of_total": 0, # filled in by caller + "affected_kernels": [ + events[i - 1]["name"], + events[i]["name"], + ], + "severity": severity, + }) + + # Compute cumulative stats + total_gap = sum(g["gap_us"] for g in gaps) + total_dur = sum( + (e["end"] - events[0]["ts"]) for e in events[-1:] + ) if events else 0 + + for g in gaps: + g["cumulative_gap_us"] = round(total_gap, 1) + g["pct_of_total"] = round(g["gap_us"] / total_dur * 100, 2) if total_dur > 0 else 0 + + return gaps + + +def build_overlap_report( + trace_data: Dict[str, Any], + batch_size: int, + stage: str, + *, + gap_threshold_us: float = 10.0, +) -> Dict[str, Any]: + """Produce the full overlap.json payload.""" + gaps = detect_gaps(trace_data, gap_threshold_us=gap_threshold_us) + total_gap = sum(g["gap_us"] for g in gaps) + total_dur = sum( + evt.get("dur", 0) for evt in + (trace_data.get("traceEvents", []) or []) + if evt.get("cat") == "kernel" + ) + + return { + "batch_size": batch_size, + "stage": stage, + "gaps": gaps, + "summary": { + "total_gap_us": round(total_gap, 1), + "total_gap_pct": round(total_gap / total_dur * 100, 2) if total_dur > 0 else 0, + "cuda_graph_effective": len(gaps) < 5, + }, + } diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/phases.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/phases.py new file mode 100644 index 00000000..a722f6a7 --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/phases.py @@ -0,0 +1,55 @@ +"""Phase graph for sglang_trace_analyze. + +Linear pipeline: MAPPING -> BENCHMARK -> ANALYZE -> HINTS -> SUMMARIZE -> done. + +The WebUI state-graph endpoint reads ``terminal_phases`` and +``graph_payload()`` from this module. +""" + +from __future__ import annotations + +from typing import Any, Dict, List, Optional + +# Ordered list of phases in the pipeline. +PHASES: List[str] = ["mapping", "benchmark", "analyze", "hints", "summarize"] + +# Phases that signal the task is done (whichever is current at exit). +TERMINAL: set[str] = {"done", "failed"} + + +def terminal_phases() -> set[str]: + return TERMINAL + + +def next_phase(current: str) -> str: + """Linear advance. Returns "done" at the end.""" + try: + idx = PHASES.index(current) + if idx + 1 < len(PHASES): + return PHASES[idx + 1] + return "done" + except ValueError: + return "done" + + +def graph_payload( + current: str = "idle", + last_outcome: Optional[str] = None, + last_label: Optional[str] = None, +) -> Dict[str, Any]: + """Return a mermaid-friendly description of the phase graph.""" + nodes = [] + edges = [] + for i, p in enumerate(PHASES): + nodes.append({"id": p, "label": p.upper()}) + if i > 0: + edges.append({"from": PHASES[i - 1], "to": p}) + edges.append({"from": PHASES[-1], "to": "done"}) + nodes.append({"id": "done", "label": "DONE"}) + return { + "nodes": nodes, + "edges": edges, + "current": current, + "last_outcome": last_outcome, + "last_transition_label": last_label, + } diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/pipeline.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/pipeline.py new file mode 100644 index 00000000..930005da --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/pipeline.py @@ -0,0 +1,564 @@ +"""Pipeline — the sglang_trace_analyze core iteration loop. + +Five-phase linear pipeline: + + MAPPING → BENCHMARK → ANALYZE → HINTS → SUMMARIZE → done + +Each phase may internally iterate (e.g. ANALYZE loops over batch_sizes × +stages). All analysis outputs are written to ``state_dir/analysis/`` as +the authoritative source of truth. +""" + +from __future__ import annotations + +import json +import re +import subprocess +import time +from pathlib import Path +from typing import Any, Dict, List, Optional + +from metainfer.orchestrator.requirements import req_field, req_field_int +from metainfer.orchestrator.state import StateStore + +from .gpu_specs import GPU_SPECS, GpuSpec +from .iteration_record import ( + AnalyzeRecord, + BenchmarkRecord, + HintsRecord, + MappingRecord, + SummarizeRecord, +) +from .phases import next_phase + + +def _load_json(path: Path, default: Any = None) -> Any: + if not path.exists(): + return default + try: + return json.loads(path.read_text(encoding="utf-8")) + except (ValueError, OSError): + return default + + +def _parse_batch_sizes(raw: str) -> List[int]: + """Parse comma-separated batch sizes, e.g. "1,4,8,16" → [1,4,8,16].""" + return [int(x.strip()) for x in raw.split(",") if x.strip()] + + +def _iter_n(store: StateStore) -> int: + """Next iteration number for timeline ordering. + + Because this pipeline runs phases sequentially with a single iteration + counter (not per-phase counters), we use a simple global counter. + """ + run = store.load_run() + return run.current_iteration + 1 if run else 1 + + +class Pipeline: + """Five-phase profiler-analysis pipeline.""" + + def __init__( + self, + req: Dict[str, Any], + store: StateStore, + state_dir: Path, + workspace_dir: Path, + ): + self.req = req + self.store = store + self.state_dir = state_dir + self.workspace_dir = workspace_dir + self._analysis_dir = state_dir / "analysis" + + # Extract form fields + self.model_path = req_field(req, "model_path", default="") + self.version = req_field(req, "version", default="dev") + self.batch_sizes = _parse_batch_sizes( + req_field(req, "batch_sizes", default="1") + ) + self.mapping_batch_size = req_field_int(req, "mapping_batch_size", default=8) + self.input_len = req_field_int(req, "input_len", default=512) + self.output_len = req_field_int(req, "output_len", default=2000) + self.tp_size = req_field_int(req, "tp_size", default=1) + self.pp_size = req_field_int(req, "pp_size", default=1) + gpu_label = req_field(req, "gpu_model", default="K100") + self.gpu_spec: GpuSpec = GPU_SPECS.get(gpu_label, GPU_SPECS["K100"]) + + # Only decode stage for now + self.stages = ["decode"] # future: ["prefill", "decode"] + + # ------------------------------------------------------------------ # + # Public entry point + # ------------------------------------------------------------------ # + + def run(self) -> None: + """Run the full pipeline.""" + run = self.store.load_run() + phase = run.current_phase or "mapping" + + while phase not in ("done", "failed"): + self.store.update_run(current_phase=phase) + self.store.append_timeline("phase_enter", {"phase": phase}) + + method = getattr(self, f"_run_{phase}", None) + if method is None: + print(f"[pipeline] unknown phase {phase!r}, stopping") + break + + try: + ok = method() + except Exception as exc: + print(f"[pipeline] phase {phase} crashed: {exc}") + self.store.append_timeline("phase_error", {"phase": phase, "error": str(exc)}) + self.store.update_run(finished=True, final_status="failed") + return + + if not ok: + print(f"[pipeline] phase {phase} returned failure, stopping") + self.store.update_run(finished=True, final_status="failed") + return + + self.store.append_timeline("phase_exit", {"phase": phase}) + phase = next_phase(phase) + + self.store.update_run(finished=True, final_status="success", + current_phase="done") + self.store.append_timeline("run_done", {"status": "success"}) + + # ================================================================== # + # Phase: MAPPING + # ================================================================== # + + def _run_mapping(self) -> bool: + """Run mapping benchmark (--disable-cuda-graph) then build the + kernel→model-structure mapping from call stacks.""" + print("[pipeline] === MAPPING phase ===") + n = _iter_n(self.store) + rec = MappingRecord(batch_size=self.mapping_batch_size) + rec.start() + self.store.write_iteration(n, rec.to_dict()) + self.store.update_run(current_iteration=n) + + trace_dir = self.workspace_dir / "traces" / "mapping" + + # 1. Generate bench_config.json + bench_config = self._build_bench_config( + batch_sizes=[self.mapping_batch_size], + output_dir=str(self.workspace_dir / "traces"), + ) + config_path = self.state_dir / "bench_config.json" + config_path.write_text(json.dumps(bench_config, indent=2)) + + # 2. Run mapping benchmark + script = Path(__file__).resolve().parent / "run_benchmark.py" + print(f"[pipeline] running mapping benchmark (batch={self.mapping_batch_size})...") + try: + subprocess.run( + [ + "python", str(script), + "--config", str(config_path), + "--mapping-only", + ], + check=True, + timeout=3600, + ) + except subprocess.TimeoutExpired: + rec.fail("mapping benchmark timed out") + self.store.write_iteration(n, rec.to_dict()) + return False + except subprocess.CalledProcessError as e: + rec.fail(f"mapping benchmark exit code {e.returncode}") + self.store.write_iteration(n, rec.to_dict()) + return False + + # 3. Parse trace → build mapping table (rule engine) + decode_trace_dir = trace_dir / "decode" + if not decode_trace_dir.exists(): + # try: the wrapper may have used --profile-by-stage naming + candidates = sorted(trace_dir.glob("*.json.gz")) + if not candidates: + rec.fail("no trace files found after mapping benchmark") + self.store.write_iteration(n, rec.to_dict()) + return False + trace_path = candidates[0] # best effort + else: + traces = sorted(decode_trace_dir.glob("*.trace.json.gz")) + if not traces: + rec.fail("no trace files in mapping/decode/") + self.store.write_iteration(n, rec.to_dict()) + return False + trace_path = traces[0] + + print(f"[pipeline] parsing trace: {trace_path}") + mapping_entries = self._build_mapping(trace_path) + if not mapping_entries: + rec.fail("mapping produced zero entries — trace may be empty or format unsupported") + self.store.write_iteration(n, rec.to_dict()) + return False + + # 4. LLM sanity check (placeholder — calls sub-agent when available) + mapping_entries = self._llm_mapping_sanity_check(mapping_entries) + + # 5. Write mapping.json + self._analysis_dir.mkdir(parents=True, exist_ok=True) + mapping_file = self._analysis_dir / "mapping.json" + mapping_file.write_text(json.dumps({ + "model": self.model_path, + "gpu": self.gpu_spec.label, + "mapping_batch_size": self.mapping_batch_size, + "entries": mapping_entries, + }, indent=2)) + + confidence_issues = sum( + 1 for e in mapping_entries + if e.get("confidence", "high") == "low" + ) + + rec.done( + trace_dir=str(trace_dir), + kernel_count=len(mapping_entries), + confidence_issues=confidence_issues, + ) + self.store.write_iteration(n, rec.to_dict()) + return True + + # ------------------------------------------------------------------ # + # Phase: BENCHMARK + # ------------------------------------------------------------------ # + + def _run_benchmark(self) -> bool: + """Run formal benchmarks (CUDA Graph ON) for each batch size.""" + print("[pipeline] === BENCHMARK phase ===") + + bench_config = self._build_bench_config( + batch_sizes=self.batch_sizes, + output_dir=str(self.workspace_dir / "traces"), + ) + config_path = self.state_dir / "bench_config.json" + config_path.write_text(json.dumps(bench_config, indent=2)) + + script = Path(__file__).resolve().parent / "run_benchmark.py" + + all_ok = True + for bs in self.batch_sizes: + n = _iter_n(self.store) + rec = BenchmarkRecord(batch_size=bs) + rec.start() + self.store.write_iteration(n, rec.to_dict()) + self.store.update_run(current_iteration=n) + + trace_dir = self.workspace_dir / "traces" / f"bs_{bs}" + print(f"[pipeline] batch_size={bs}") + + try: + subprocess.run( + [ + "python", str(script), + "--config", str(config_path), + "--formal-only", + "--single-batch", str(bs), + ], + check=True, + timeout=3600, + ) + except subprocess.TimeoutExpired: + rec.fail("timed out") + self.store.write_iteration(n, rec.to_dict()) + all_ok = False + continue + except subprocess.CalledProcessError as e: + rec.fail(f"exit code {e.returncode}") + self.store.write_iteration(n, rec.to_dict()) + all_ok = False + continue + + # Extract throughput/latency from sglang output log if available + rec.done(trace_dir=str(trace_dir)) + self.store.write_iteration(n, rec.to_dict()) + + # We continue even if some batches failed — ANALYZE skips them + return all_ok or any( + (self.workspace_dir / "traces" / f"bs_{bs}" / "decode").exists() + for bs in self.batch_sizes + ) + + # ================================================================== # + # Phase: ANALYZE + # ================================================================== # + + def _run_analyze(self) -> bool: + """Analyze each (batch_size, stage) pair that has a trace.""" + print("[pipeline] === ANALYZE phase ===") + + mapping = _load_json(self._analysis_dir / "mapping.json", {}) + mapping_entries = mapping.get("entries", []) + if not mapping_entries: + print("[pipeline] WARNING: no mapping entries — analysis may be incomplete") + + any_ok = False + for bs in self.batch_sizes: + for stage in self.stages: + trace_dir = self.workspace_dir / "traces" / f"bs_{bs}" / stage + if not trace_dir.exists(): + print(f"[pipeline] skipping bs_{bs}/{stage} — no trace dir") + continue + + traces = sorted(trace_dir.glob("*.trace.json.gz")) + if not traces: + print(f"[pipeline] skipping bs_{bs}/{stage} — no trace files") + continue + + n = _iter_n(self.store) + rec = AnalyzeRecord(batch_size=bs, stage=stage) + rec.start() + self.store.write_iteration(n, rec.to_dict()) + self.store.update_run(current_iteration=n) + + print(f"[pipeline] analyzing bs_{bs}/{stage} ({traces[0].name})") + try: + result = self._analyze_one( + traces[0], mapping_entries, bs, stage + ) + except Exception as exc: + rec.fail(str(exc)) + self.store.write_iteration(n, rec.to_dict()) + continue + + # Write output files + out_dir = self._analysis_dir / "batches" / f"bs_{bs}" / stage + out_dir.mkdir(parents=True, exist_ok=True) + + (out_dir / "kernel_table.json").write_text( + json.dumps(result["kernel_table"], indent=2)) + (out_dir / "overlap.json").write_text( + json.dumps(result["overlap"], indent=2)) + (out_dir / "fuse.json").write_text( + json.dumps(result["fuse"], indent=2)) + + mfu_vals = [ + k.get("mfu", 0) for k in result["kernel_table"].get("kernels", []) + if k.get("mfu") is not None + ] + top_kernels = result["kernel_table"].get("kernels", []) + rec.done( + kernel_count=len(top_kernels), + top_kernel=top_kernels[0]["kernel_name"] if top_kernels else None, + top_kernel_pct=top_kernels[0]["time_pct"] if top_kernels else None, + mfu_avg=round(sum(mfu_vals) / len(mfu_vals), 1) if mfu_vals else None, + fuse_hits=len(result["fuse"].get("matches", [])), + ) + self.store.write_iteration(n, rec.to_dict()) + any_ok = True + + return any_ok + + # ================================================================== # + # Phase: HINTS + # ================================================================== # + + def _run_hints(self) -> bool: + """Generate LLM optimization hints from all analysis results.""" + print("[pipeline] === HINTS phase ===") + n = _iter_n(self.store) + rec = HintsRecord( + model_used=self.model_path, + batch_count=len(self.batch_sizes), + ) + rec.start() + self.store.write_iteration(n, rec.to_dict()) + self.store.update_run(current_iteration=n) + + # Collect summaries from all analyzed batches + kernel_summaries = [] + overlap_summaries = [] + fuse_summaries = [] + + for bs in self.batch_sizes: + for stage in self.stages: + out_dir = self._analysis_dir / "batches" / f"bs_{bs}" / stage + kt = _load_json(out_dir / "kernel_table.json") + ov = _load_json(out_dir / "overlap.json") + fu = _load_json(out_dir / "fuse.json") + if kt: + top3 = (kt.get("kernels", []) or [])[:3] + kernel_summaries.append({ + "batch_size": bs, "stage": stage, + "top_kernels": top3, + }) + if ov: + overlap_summaries.append(ov) + if fu: + fuse_summaries.append(fu) + + if not kernel_summaries: + print("[pipeline] no kernel tables — skipping hints") + rec.fail("no analysis data available") + self.store.write_iteration(n, rec.to_dict()) + return True # not fatal — hints are optional + + # Generate hints (placeholder — real impl calls LLM sub-agent) + hints = self._llm_generate_hints( + kernel_summaries, overlap_summaries, fuse_summaries + ) + self._analysis_dir.mkdir(parents=True, exist_ok=True) + (self._analysis_dir / "hints.json").write_text( + json.dumps(hints, indent=2)) + + rec.done() + self.store.write_iteration(n, rec.to_dict()) + return True + + # ================================================================== # + # Phase: SUMMARIZE + # ================================================================== # + + def _run_summarize(self) -> bool: + """Aggregate cross-batch summary.""" + print("[pipeline] === SUMMARIZE phase ===") + n = _iter_n(self.store) + rec = SummarizeRecord(batch_count=len(self.batch_sizes)) + rec.start() + self.store.write_iteration(n, rec.to_dict()) + self.store.update_run(current_iteration=n) + + batch_summaries = [] + best_batch = None + best_mfu = None + + for bs in self.batch_sizes: + for stage in ["decode"]: + out_dir = self._analysis_dir / "batches" / f"bs_{bs}" / stage + kt = _load_json(out_dir / "kernel_table.json") + if not kt: + batch_summaries.append({ + "batch_size": bs, "stage": stage, + "status": "missing", + }) + continue + + kernels = kt.get("kernels", []) or [] + mfu_vals = [k.get("mfu", 0) for k in kernels if k.get("mfu")] + avg_mfu = round(sum(mfu_vals) / len(mfu_vals), 1) if mfu_vals else None + top = kernels[0] if kernels else {} + + info = { + "batch_size": bs, + "stage": stage, + "top_kernel": top.get("kernel_name"), + "top_kernel_pct": top.get("time_pct"), + "mfu_avg": avg_mfu, + "kernel_count": len(kernels), + } + batch_summaries.append(info) + + if avg_mfu is not None and (best_mfu is None or avg_mfu > best_mfu): + best_mfu = avg_mfu + best_batch = bs + + self._analysis_dir.mkdir(parents=True, exist_ok=True) + (self._analysis_dir / "summary.json").write_text(json.dumps({ + "model": self.model_path, + "gpu": self.gpu_spec.label, + "batches": batch_summaries, + }, indent=2)) + + rec.done(best_batch=best_batch, best_mfu=best_mfu) + self.store.write_iteration(n, rec.to_dict()) + return True + + # ================================================================== # + # Helpers + # ================================================================== # + + def _build_bench_config( + self, batch_sizes: List[int], output_dir: str, + ) -> Dict[str, Any]: + return { + "model_path": self.model_path, + "version": self.version, + "batch_sizes": batch_sizes, + "mapping_batch_size": self.mapping_batch_size, + "input_len": self.input_len, + "output_len": self.output_len, + "tp_size": self.tp_size, + "pp_size": self.pp_size, + "output_dir": output_dir, + } + + def _build_mapping(self, trace_path: Path) -> List[Dict[str, Any]]: + """Parse a torch profiler Chrome trace and extract kernel→layer + mappings from call stacks. + + Placeholder implementation — real logic will live in + ``trace_parser.py`` and ``structure_mapper.py``. + """ + # TODO: implement trace_parser.py + print("[pipeline] _build_mapping: parsing trace (placeholder)") + return [] + + def _llm_mapping_sanity_check( + self, entries: List[Dict[str, Any]] + ) -> List[Dict[str, Any]]: + """Run LLM sanity check on mapping entries. + + Placeholder — real impl calls SubAgentManager. + """ + # TODO: wire SubAgentManager + for e in entries: + e.setdefault("confidence", "high") + return entries + + def _analyze_one( + self, + trace_path: Path, + mapping_entries: List[Dict[str, Any]], + bs: int, + stage: str, + ) -> Dict[str, Any]: + """Analyze a single trace file and return kernel_table, overlap, + and fuse results. + + Placeholder — real logic in trace_parser / flops_calculator / + overlap_detector / fuse_matcher. + """ + # TODO: implement real analysis pipeline + return { + "kernel_table": { + "model": self.model_path, + "gpu": self.gpu_spec.label, + "batch_size": bs, + "stage": stage, + "kernels": [], + }, + "overlap": { + "batch_size": bs, + "stage": stage, + "gaps": [], + "summary": {"total_gap_us": 0, "total_gap_pct": 0, "cuda_graph_effective": True}, + }, + "fuse": { + "batch_size": bs, + "stage": stage, + "matches": [], + }, + } + + def _llm_generate_hints( + self, + kernel_summaries: list, + overlap_summaries: list, + fuse_summaries: list, + ) -> Dict[str, Any]: + """Generate optimization hints via LLM. + + Placeholder — real impl calls SubAgentManager. + """ + return { + "bottleneck": {"kernel_or_pattern": "TBD", "reason": "", "impact_pct": 0}, + "suggestions": [], + "surprises": [], + "status": "skipped", + "reason": "LLM hints not yet wired", + } diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/plugin.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/plugin.py new file mode 100644 index 00000000..d3d14812 --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/plugin.py @@ -0,0 +1,10 @@ +"""TaskPlugin descriptor for sglang_trace_analyze.""" + +from metainfer.orchestrator.tasks.base import TaskPlugin + +PLUGIN = TaskPlugin( + task_type="sglang_trace_analyze", + cli_module="metainfer.tasks.sglang_trace_analyze.orchestrator.cli", + phases_module="metainfer.tasks.sglang_trace_analyze.orchestrator.phases", + diagnostic_globs=("*",), +) diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/prompts.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/prompts.py new file mode 100644 index 00000000..2b4994b4 --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/prompts.py @@ -0,0 +1,117 @@ +"""LLM prompts for sglang_trace_analyze. + +Two prompt families: +1. **mapping_sanity_check** — validate the kernel-to-model-structure mapping. +2. **optimization_hints** — generate actionable optimization suggestions from + the full analysis (kernel tables + overlap + fuse results). +""" + +from __future__ import annotations + + +def mapping_sanity_check_prompt( + mapping_json: str, + model_config_json: str, + gpu_label: str, +) -> str: + """Prompt for LLM to sanity-check a kernel → model-structure mapping.""" + return f"""You are a GPU inference optimization expert. Review the following +kernel-to-model-structure mapping that was auto-generated from a torch +profiler trace's call stacks. + +## Model config.json +```json +{model_config_json} +``` + +## Auto-generated mapping (excerpt — full file too large, this is the first +200 entries sorted by GPU time) +```json +{mapping_json} +``` + +## GPU +{gpu_label} + +## Tasks +1. For each mapping entry, rate its confidence: "high" (call stack clearly + points to a known layer/op), "medium" (plausible but ambiguous), or + "low" (likely wrong — kernel name and call stack don't match expected + pattern). If you're uncertain about a model architecture detail, search + the web for the model's architecture documentation before rating. +2. Flag any kernel that appears to be mapped to the wrong layer type + (e.g. a MoE kernel mapped to a dense layer, or an attention kernel + mapped to an FFN layer). +3. Flag missing mappings — kernel names that appear in the trace but have + no clear model-layer assignment. +4. Return a JSON object with this schema: + {{ + "entries": [ + {{ + "kernel_name": "...", + "confidence": "high|medium|low", + "issues": ["..."] // empty list if none + }} + ], + "summary": {{ + "high_count": N, + "medium_count": N, + "low_count": N, + "overall_assessment": "..." + }} + }} +""" + + +def optimization_hints_prompt( + kernel_tables_summary: str, + overlap_summary: str, + fuse_summary: str, + gpu_label: str, + model_name: str, +) -> str: + """Prompt for LLM to generate optimization hints from analysis results.""" + return f"""You are a GPU inference optimization expert. Review the profiling +analysis below and generate actionable optimization suggestions. + +## Model +{model_name} + +## GPU +{gpu_label} + +## Kernel Hotspot Summary (top kernels by GPU time across all batch sizes) +{kernel_tables_summary} + +## Overlap Analysis +{overlap_summary} + +## Fuse Pattern Matches +{fuse_summary} + +## Tasks +1. Identify the single biggest bottleneck and explain why it dominates. +2. List 3-5 concrete optimization directions, ordered by estimated impact. + For each: what to change, why it helps, and estimated saving (%). +3. Note any surprising or counter-intuitive findings (e.g. a kernel that + should be fast but is unexpectedly slow). +4. Return a JSON object with this schema: + {{ + "bottleneck": {{ + "kernel_or_pattern": "...", + "reason": "...", + "impact_pct": N + }}, + "suggestions": [ + {{ + "title": "...", + "what_to_change": "...", + "why": "...", + "estimated_saving_pct": N, + "difficulty": "low|medium|high", + "category": "fuse|overlap|kernel_replace|config_tune|other" + }} + ], + "surprises": ["..."] + }} +""" diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/run_benchmark.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/run_benchmark.py new file mode 100644 index 00000000..469cf07e --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/run_benchmark.py @@ -0,0 +1,152 @@ +#!/usr/bin/env python3 +"""Wrapper script for sglang.bench_one_batch_server. + +Called by the orchestrator pipeline in two modes: + + # Mapping run — one batch size, --disable-cuda-graph + python run_benchmark.py --config bench_config.json --mapping-only + + # Formal runs — one or all batch sizes, CUDA Graph ON + python run_benchmark.py --config bench_config.json --formal-only [--single-batch N] + +The benchmark is a synchronous, blocking call — the caller waits for +all batch sizes to complete. +""" + +from __future__ import annotations + +import argparse +import json +import os +import subprocess +import sys +from pathlib import Path +from typing import Dict, Any, List + + +def build_dir_name(args: Dict[str, Any]) -> str: + """Build sglang-style directory name from config.""" + parts = [args["version"], f"tp{args['tp_size']}", f"pp{args['pp_size']}"] + parts.append("graph") + return "_".join(parts) + + +def run_benchmark( + args: Dict[str, Any], + dir_name: str, + batch_size: int, + *, + disable_cuda_graph: bool = False, +) -> bool: + """Run a single bench_one_batch_server invocation.""" + output_dir = os.path.join( + args["output_dir"], "mapping" if disable_cuda_graph else f"bs_{batch_size}" + ) + profile_prefix = f"{dir_name}_" + + cmd = [ + sys.executable, "-m", "sglang.bench_one_batch_server", + "--model-path", args["model_path"], + "--tp-size", str(args["tp_size"]), + "--pp-size", str(args["pp_size"]), + "--batch-size", str(batch_size), + "--input-len", str(args["input_len"]), + "--output-len", str(args["output_len"]), + "--run-name", dir_name, + "--show-report", + "--dataset-name", "random-ids", + "--fake-prefill", + "--profile", + "--profile-start-step", "500", + "--profile-steps", "50", + "--profile-by-stage", + "--profile-prefix", profile_prefix, + "--profile-output-dir", output_dir, + "--disable-radix-cache", + "--chunked-prefill-size", "4096", + "--kv-cache-dtype", "auto", + "--disable-flashinfer-autotune", + "--enable-metrics", + ] + + if disable_cuda_graph: + cmd.append("--disable-cuda-graph") + else: + cmd.extend(["--cuda-graph-bs", str(batch_size)]) + + print(f"\n{'='*80}") + print(f"Running batch_size={batch_size}" + f"{' (CUDA Graph OFF)' if disable_cuda_graph else ''}") + print(f" profile-output-dir: {output_dir}") + print(f" profile-prefix: {profile_prefix}") + print(f"{'='*80}\n") + + try: + subprocess.run(cmd, check=True, timeout=3600) + except subprocess.TimeoutExpired: + print(f"\n[FAILED] batch_size={batch_size}: timed out after 1 hour\n") + return False + except subprocess.CalledProcessError as e: + print(f"\n[FAILED] batch_size={batch_size}: exit code {e.returncode}\n") + return False + return True + + +def main(): + parser = argparse.ArgumentParser( + description="Run sglang bench_one_batch_server with torch profiler" + ) + parser.add_argument("--config", required=True, + help="Path to JSON benchmark config") + parser.add_argument("--mapping-only", action="store_true", + help="Run only the mapping benchmark (--disable-cuda-graph, one batch)") + parser.add_argument("--formal-only", action="store_true", + help="Run formal benchmarks (CUDA Graph ON, one or all batches)") + parser.add_argument("--single-batch", type=int, default=None, + help="When --formal-only, run only this batch size") + + args = parser.parse_args() + config_path = Path(args.config) + if not config_path.exists(): + print(f"ERROR: config file not found: {args.config}") + return 1 + + with open(config_path) as f: + cfg = json.load(f) + + dir_name = build_dir_name(cfg) + + if args.mapping_only: + bs = cfg.get("mapping_batch_size", 8) + ok = run_benchmark(cfg, dir_name, bs, disable_cuda_graph=True) + return 0 if ok else 1 + + if args.formal_only: + batch_sizes: List[int] = cfg.get("batch_sizes", [1]) + if args.single_batch is not None: + if args.single_batch in batch_sizes: + batch_sizes = [args.single_batch] + else: + print(f"ERROR: --single-batch {args.single_batch} not in " + f"configured batch_sizes {batch_sizes}") + return 1 + + succeeded, failed = [], [] + for bs in batch_sizes: + ok = run_benchmark(cfg, dir_name, bs) + (succeeded if ok else failed).append(bs) + + print(f"\n{'='*80}") + print(f"Completed: {len(succeeded)} succeeded, {len(failed)} failed") + if succeeded: + print(f" Succeeded batches: {succeeded}") + if failed: + print(f" Failed batches: {failed}") + return 0 if not failed else 1 + + print("ERROR: must specify --mapping-only or --formal-only") + return 1 + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/structure_mapper.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/structure_mapper.py new file mode 100644 index 00000000..3d5c5dcb --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/structure_mapper.py @@ -0,0 +1,158 @@ +"""Map kernel names (via call stacks) to model structural elements. + +Takes the aggregated kernel list from :mod:`trace_parser` and the model's +``config.json``, then assigns each kernel to: +- ``model_layer`` — e.g. ``layer_{2..58}/attn/qkv_proj`` +- ``op_type`` — GEMM / Attention / Norm / ElementWise / MoE / NCCL / ... +- ``category`` — for grouping (MLA, MoE, GEMM, NCCL, etc.) + +Mapping is done by parsing the Python source location from the call stack +and matching it against known sglang layer source patterns. +""" + +from __future__ import annotations + +from typing import Any, Dict, List, Optional + + +def build_mapping( + kernels: List[Dict[str, Any]], + config: Dict[str, Any], +) -> List[Dict[str, Any]]: + """Build the kernel-to-model-structure mapping. + + Args: + kernels: Aggregated kernel list from ``aggregate_kernels()`` with + ``include_call_stack=True``. + config: Model ``config.json`` as a dict. + + Returns: + List of mapping entries with ``kernel_name``, ``model_layer``, + ``op_type``, ``category``, ``call_stack``, ``confidence``. + """ + entries = [] + for k in kernels: + call_stack = k.get("call_stack", "") + entry = _map_one(k["kernel_name"], call_stack, config) + entries.append(entry) + return entries + + +# ------------------------------------------------------------------ # +# Internal: pattern-based mapping +# ------------------------------------------------------------------ # + +def _map_one( + kernel_name: str, + call_stack: str, + config: Dict[str, Any], +) -> Dict[str, Any]: + """Map a single kernel to a model layer by inspecting its call stack.""" + layer = _infer_layer(call_stack, kernel_name, config) + op_type = _infer_op_type(kernel_name, call_stack) + + confidence = "high" + if not call_stack: + confidence = "low" + elif layer is None: + confidence = "medium" + + return { + "kernel_name": kernel_name, + "model_layer": layer, + "op_type": op_type, + "category": _op_type_to_category(op_type), + "call_stack": call_stack, + "confidence": confidence, + } + + +def _infer_layer( + call_stack: str, + kernel_name: str, + config: Dict[str, Any], +) -> Optional[str]: + """Extract layer information from the call stack. + + Looks for patterns like: + - ``sglang/srt/layers/...`` + - ``layer_forward`` + - ``model.py``, ``decoder.py``, ``encoder.py`` + - Module names like ``model.layers.5.self_attn`` + + Returns ``None`` if no layer info can be inferred. + """ + if not call_stack: + return None + + # Heuristic: look for sglang/srt/layers or model.layers.N patterns + lines = call_stack.strip().split("\n") + + # Pattern 1: model.layers.N in the call stack + import re + layer_pat = re.compile(r"model\.layers\.(\d+)") + # Pattern 2: sglang source files under layers/ + sglang_layer_pat = re.compile( + r"sglang/srt/layers/(attn|moe|mla|linear|norm|embed|sampler|router)" + ) + + for line in lines: + m = layer_pat.search(line) + if m: + return f"layer_{m.group(1)}" + m = sglang_layer_pat.search(line) + if m: + return f"layers/{m.group(1)}" + + # Fallback: use kernel name heuristics + if "attn" in kernel_name.lower() or "attention" in kernel_name.lower(): + return "attention (unknown layer)" + if "moe" in kernel_name.lower(): + return "moe (unknown layer)" + if "gemm" in kernel_name.lower() or "linear" in kernel_name.lower(): + return "linear (unknown layer)" + + return None + + +def _infer_op_type(kernel_name: str, call_stack: str) -> str: + """Infer the op type from kernel name and call stack.""" + name_lower = kernel_name.lower() + if any(k in name_lower for k in ("attn", "attention", "flash_fwd", "flash_attn")): + return "Attention" + if any(k in name_lower for k in ("moe", "fused_moe")): + return "MoE" + if any(k in name_lower for k in ("gemm", "linear", "matmul", "w8a8", "fp8")): + return "GEMM" + if any(k in name_lower for k in ("rms", "norm", "layernorm", "layer_norm")): + return "Norm" + if any(k in name_lower for k in ("nccl", "allreduce", "allgather", "broadcast")): + return "NCCL" + if any(k in name_lower for k in ("hadamard", "rotate", "rope")): + return "Transform" + if any(k in name_lower for k in ("copy", "memcpy", "memset")): + return "Memory" + if any(k in name_lower for k in ("silu", "gelu", "swiglu", "activation", "act_and_mul")): + return "Activation" + if any(k in name_lower for k in ("topk", "top_k", "index", "gather", "scatter", "sort")): + return "Indexing" + if any(k in name_lower for k in ("quant", "dequant", "fp8_scale")): + return "Quantization" + return "Other" + + +def _op_type_to_category(op_type: str) -> str: + """Map an op_type to a display category.""" + mapping = { + "Attention": "Attention", + "MoE": "MoE", + "GEMM": "GEMM", + "Norm": "Norm", + "NCCL": "NCCL", + "Transform": "Transform", + "Memory": "Memory", + "Activation": "Activation", + "Indexing": "Indexing", + "Quantization": "Quantization", + } + return mapping.get(op_type, "Other") diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/trace_parser.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/trace_parser.py new file mode 100644 index 00000000..83bc8a58 --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/trace_parser.py @@ -0,0 +1,140 @@ +"""Chrome trace JSON parser + kernel aggregation. + +Loads a ``torch.profiler`` Chrome trace (``.json`` or ``.json.gz``) and +produces an aggregated kernel table: one row per unique kernel name, +sorted by total GPU duration descending. + +In the MAPPING phase this also extracts call-stack information for +structure mapping. In the ANALYZE phase it aggregates CUDA Graph replay +events into per-kernel durations. +""" + +from __future__ import annotations + +import gzip +import json +from collections import defaultdict +from pathlib import Path +from typing import Any, Dict, List, Optional + + +def _open_trace(trace_path: Path): + """Open a trace file — transparently handles .gz compression.""" + if trace_path.suffix == ".gz": + return gzip.open(trace_path, "rt", encoding="utf-8") + return open(trace_path, "r", encoding="utf-8") + + +def parse_trace(trace_path: Path) -> Dict[str, Any]: + """Load a Chrome trace JSON and return the top-level document. + + Returns: + Dict with keys: ``traceEvents``, ``displayTimeUnit``, etc. + """ + with _open_trace(trace_path) as f: + data = json.load(f) + return data + + +def aggregate_kernels( + trace_data: Dict[str, Any], + *, + include_call_stack: bool = False, +) -> List[Dict[str, Any]]: + """Aggregate GPU kernel events by kernel name. + + Args: + trace_data: Parsed Chrome trace JSON. + include_call_stack: If True, preserve ``call_stack`` from the first + occurrence of each unique kernel name. + + Returns: + List of kernel dicts sorted by ``total_dur_us`` descending. Each dict: + ``kernel_name``, ``total_dur_us``, ``count``, ``call_stack`` (optional). + """ + trace_events = trace_data.get("traceEvents", []) + if not trace_events: + # sglang sometimes wraps in a list directly + if isinstance(trace_data, list): + trace_events = trace_data + else: + return [] + + # Filter GPU kernel events + kernels: Dict[str, Dict[str, Any]] = {} + for evt in trace_events: + cat = evt.get("cat", "") + name = evt.get("name", "") + dur = evt.get("dur", 0) + + # Torch profiler GPU kernel events: cat="kernel", name like + # "triton_fused_moe_kernel" or "void at::native::..." + if cat != "kernel" or dur <= 0: + continue + + if name not in kernels: + entry: Dict[str, Any] = { + "kernel_name": name, + "total_dur_us": 0, + "count": 0, + } + if include_call_stack: + args = evt.get("args", {}) or {} + call_stack = args.get("call stack", "") + if call_stack: + entry["call_stack"] = call_stack + kernels[name] = entry + + kernels[name]["total_dur_us"] += dur + kernels[name]["count"] += 1 + + # Sort by total duration descending + result = sorted( + kernels.values(), key=lambda k: k["total_dur_us"], reverse=True + ) + return result + + +def aggregate_kernels_with_dims( + trace_data: Dict[str, Any], +) -> List[Dict[str, Any]]: + """Like :func:`aggregate_kernels`, but also collects Input Dims from + ``args["Input Dims"]`` for shape-aware kernels (GEMM, attention). + + This is only meaningful when the trace was captured WITHOUT CUDA Graph + (i.e. during the MAPPING phase), because CUDA Graph replay hides + individual kernel dims. + """ + trace_events = trace_data.get("traceEvents", []) + if isinstance(trace_data, list): + trace_events = trace_data + + kernels: Dict[str, Dict[str, Any]] = {} + for evt in trace_events: + cat = evt.get("cat", "") + name = evt.get("name", "") + dur = evt.get("dur", 0) + if cat != "kernel" or dur <= 0: + continue + + if name not in kernels: + args = evt.get("args", {}) or {} + entry: Dict[str, Any] = { + "kernel_name": name, + "total_dur_us": 0, + "count": 0, + "input_dims": [], + "call_stack": args.get("call stack", ""), + } + kernels[name] = entry + + kernels[name]["total_dur_us"] += dur + kernels[name]["count"] += 1 + args = evt.get("args", {}) or {} + dims = args.get("Input Dims", []) + if dims and dims not in kernels[name]["input_dims"]: + kernels[name]["input_dims"].append(dims) + + return sorted( + kernels.values(), key=lambda k: k["total_dur_us"], reverse=True + ) diff --git a/metainfer/tasks/sglang_trace_analyze/server/__init__.py b/metainfer/tasks/sglang_trace_analyze/server/__init__.py new file mode 100644 index 00000000..3da30f91 --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/server/__init__.py @@ -0,0 +1 @@ +"""Server-side plguin for sglang_trace_analyze.""" diff --git a/metainfer/tasks/sglang_trace_analyze/server/_state_readers.py b/metainfer/tasks/sglang_trace_analyze/server/_state_readers.py new file mode 100644 index 00000000..c7e5bb45 --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/server/_state_readers.py @@ -0,0 +1,51 @@ +"""State-dir readers for sglang_trace_analyze. + +Reads the authoritative analysis JSON files from +``/analysis/``. +""" + +from __future__ import annotations + +import json +from pathlib import Path +from typing import Any, Dict, Optional + + +def _load_json(path: Path) -> Optional[Any]: + if not path.exists(): + return None + try: + return json.loads(path.read_text(encoding="utf-8")) + except (ValueError, OSError): + return None + + +def read_summary(state_dir: Path) -> Optional[Dict[str, Any]]: + return _load_json(state_dir / "analysis" / "summary.json") + + +def read_mapping(state_dir: Path) -> Optional[Dict[str, Any]]: + return _load_json(state_dir / "analysis" / "mapping.json") + + +def read_hints(state_dir: Path) -> Optional[Dict[str, Any]]: + return _load_json(state_dir / "analysis" / "hints.json") + + +def read_batch_detail( + state_dir: Path, bs: int, stage: str +) -> Optional[Dict[str, Any]]: + """Return the combined kernel_table + overlap + fuse for one + (batch_size, stage) pair. + """ + base = state_dir / "analysis" / "batches" / f"bs_{bs}" / stage + kernel_table = _load_json(base / "kernel_table.json") + overlap = _load_json(base / "overlap.json") + fuse = _load_json(base / "fuse.json") + if kernel_table is None and overlap is None and fuse is None: + return None + return { + "kernel_table": kernel_table, + "overlap": overlap, + "fuse": fuse, + } diff --git a/metainfer/tasks/sglang_trace_analyze/server/plugin.py b/metainfer/tasks/sglang_trace_analyze/server/plugin.py new file mode 100644 index 00000000..6fdaa277 --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/server/plugin.py @@ -0,0 +1,31 @@ +"""WebPlugin for sglang_trace_analyze — registers routes + detail view.""" + +from __future__ import annotations + +from pathlib import Path + +from metainfer.server.registry import WebPlugin, register + +from .routes import build_router + +PLUGIN_TYPE = "sglang_trace_analyze" +_FRONTEND_DIR = Path(__file__).resolve().parent.parent / "static" +_STATIC_PREFIX = f"/static/plugins/{PLUGIN_TYPE}" + +_IMPORTMAP_ENTRIES: dict = {} + +plugin = WebPlugin( + type=PLUGIN_TYPE, + label="SGLang Trace Analyze", + description=( + "Profile a model with SGLang's torch profiler across multiple batch " + "sizes, then analyze kernel hotspots, TFLOPS/MFU, operator-to-model-" + "structure mapping, fuse opportunities, and generate LLM-powered " + "optimization hints." + ), + build_router=build_router, + frontend_dir=_FRONTEND_DIR, + importmap_entries=_IMPORTMAP_ENTRIES, +) + +register(plugin) diff --git a/metainfer/tasks/sglang_trace_analyze/server/routes.py b/metainfer/tasks/sglang_trace_analyze/server/routes.py new file mode 100644 index 00000000..3bf73822 --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/server/routes.py @@ -0,0 +1,68 @@ +"""FastAPI router for sglang_trace_analyze. + +Routes mounted under ``/api/sglang_trace_analyze/{task_id}``: + + GET /summary → summary.json + GET /mapping → mapping.json + GET /hints → hints.json + GET /batch/{bs}/{stage} → {kernel_table, overlap, fuse} +""" + +from __future__ import annotations + +from fastapi import APIRouter, HTTPException + +from metainfer.server._helpers import ( + require_task_type, + state_dir_for, + task_or_404, +) +from . import _state_readers + +PLUGIN_TYPE = "sglang_trace_analyze" + + +def build_router(plugin) -> APIRouter: + router = APIRouter() + + @router.get("/summary") + def get_summary(task_id: str): + entry = task_or_404(task_id) + require_task_type(entry, PLUGIN_TYPE) + data = _state_readers.read_summary(state_dir_for(entry)) + if data is None: + raise HTTPException(404, "summary not yet available") + return data + + @router.get("/mapping") + def get_mapping(task_id: str): + entry = task_or_404(task_id) + require_task_type(entry, PLUGIN_TYPE) + data = _state_readers.read_mapping(state_dir_for(entry)) + if data is None: + raise HTTPException(404, "mapping not yet available") + return data + + @router.get("/hints") + def get_hints(task_id: str): + entry = task_or_404(task_id) + require_task_type(entry, PLUGIN_TYPE) + data = _state_readers.read_hints(state_dir_for(entry)) + if data is None: + raise HTTPException(404, "hints not yet available") + return data + + @router.get("/batch/{bs}/{stage}") + def get_batch_detail(task_id: str, bs: int, stage: str): + entry = task_or_404(task_id) + require_task_type(entry, PLUGIN_TYPE) + data = _state_readers.read_batch_detail( + state_dir_for(entry), bs, stage + ) + if data is None: + raise HTTPException( + 404, f"no analysis data for batch {bs}/{stage}" + ) + return data + + return router diff --git a/metainfer/tasks/sglang_trace_analyze/tests/__init__.py b/metainfer/tasks/sglang_trace_analyze/tests/__init__.py new file mode 100644 index 00000000..e69de29b diff --git a/metainfer/tasks/sglang_trace_analyze/tests/test_flops_calculator.py b/metainfer/tasks/sglang_trace_analyze/tests/test_flops_calculator.py new file mode 100644 index 00000000..ef31820b --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/tests/test_flops_calculator.py @@ -0,0 +1,78 @@ +"""FLOPs calculator tests.""" + +from ..orchestrator.gpu_specs import GpuSpec +from ..orchestrator.flops_calculator import ( + _estimate_flops, + _estimate_bytes, + calculate_mfu, +) + + +K100 = GpuSpec( + label="K100", + fp32_tflops=49, + tf32_tflops=98, + bf16_tflops=192, + fp16_tflops=192, + int8_tops=392, + bandwidth_gb_s=700, +) + + +def test_estimate_flops_gemm_3d(): + # M=4096, K=2048, N=512 → 2*4096*2048*512 = 8,589,934,592 + flops = _estimate_flops("GEMM", [[4096, 2048, 512]], batch_size=1) + assert flops == 2 * 4096 * 2048 * 512 + + +def test_estimate_flops_gemm_batched(): + # B=4, M=1024, N=512, K=2048 → 2*4*1024*2048*512 + flops = _estimate_flops("GEMM", [[4, 1024, 512, 2048]], batch_size=4) + assert flops == 2 * 4 * 1024 * 2048 * 512 + + +def test_estimate_flops_no_dims(): + assert _estimate_flops("GEMM", [], batch_size=8) == 0 + + +def test_estimate_bytes_gemm(): + bytes_moved = _estimate_bytes("GEMM", [[4096, 2048, 512]], batch_size=1) + # (4096*2048 + 2048*512 + 4096*512) * 2 bytes + expected = (4096 * 2048 + 2048 * 512 + 4096 * 512) * 2 + assert bytes_moved == expected + + +def test_calculate_mfu_basic(): + # 2*4096*2048*512 = 8.59e9 FLOPs. At 10 us this is ~859 TFLOPS + # (far above K100 peak), but this is synthetic — we just verify + # the fields are populated and reasonable. + kernels = [ + { + "kernel_name": "triton_gemm", + "total_dur_us": 50, # 50 us for 8.6e9 FLOPs = 172 TFLOPS + "count": 1, + "input_dims": [[4096, 2048, 512]], + "op_type": "GEMM", + } + ] + result = calculate_mfu(kernels, K100, batch_size=1, dtype="bf16") + k = result[0] + assert k["tflops_theoretical"] == 192 + assert k["tflops_actual"] > 0 + assert k["mfu"] > 0 + assert k["bound"] in ("compute", "memory") + + +def test_calculate_mfu_no_dims(): + kernels = [ + { + "kernel_name": "cuda_graph_replay", + "total_dur_us": 500_000, + "count": 1, + "input_dims": [], + "op_type": "Other", + } + ] + result = calculate_mfu(kernels, K100, batch_size=8, dtype="bf16") + assert result[0]["tflops_actual"] == 0 + assert result[0]["mfu"] == 0 diff --git a/metainfer/tasks/sglang_trace_analyze/tests/test_fuse_matcher.py b/metainfer/tasks/sglang_trace_analyze/tests/test_fuse_matcher.py new file mode 100644 index 00000000..47fcd8f1 --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/tests/test_fuse_matcher.py @@ -0,0 +1,37 @@ +"""Fuse matcher tests.""" + +from ..orchestrator.fuse_matcher import _match_consecutive, match_fuse_patterns + + +def test_match_consecutive_found(): + names = ["abc", "rms_norm", "triton_gemm", "add"] + pattern = ["rms_norm", "gemm"] + result = _match_consecutive(names, pattern) + assert result == ["rms_norm", "triton_gemm"] + + +def test_match_consecutive_not_found(): + names = ["abc", "rms_norm", "add"] + pattern = ["rms_norm", "gemm"] + result = _match_consecutive(names, pattern) + assert result == [] + + +def test_match_consecutive_short_list(): + names = ["abc"] + pattern = ["a", "b"] + result = _match_consecutive(names, pattern) + assert result == [] + + +def test_match_fuse_patterns_with_known_kernels(): + kernels = [ + {"kernel_name": "abc"}, + {"kernel_name": "triton_gemm"}, + {"kernel_name": "ncclAllReduce"}, + {"kernel_name": "triton_gemm"}, + ] + matches = match_fuse_patterns(kernels) + # The "nccl_allreduce + gemm (no overlap)" pattern should fire + pattern_names = [m["pattern"] for m in matches] + assert "nccl_allreduce + gemm (no overlap)" in pattern_names diff --git a/metainfer/tasks/sglang_trace_analyze/tests/test_gpu_specs.py b/metainfer/tasks/sglang_trace_analyze/tests/test_gpu_specs.py new file mode 100644 index 00000000..0c8cb5b0 --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/tests/test_gpu_specs.py @@ -0,0 +1,17 @@ +"""GPU specs lookup tests.""" + +from ..orchestrator.gpu_specs import GPU_SPECS, GpuSpec + + +def test_gpu_specs_known(): + for label in ("K100", "A100_80G", "H100", "B200"): + spec = GPU_SPECS.get(label) + assert spec is not None, f"missing spec for {label}" + assert spec.bf16_tflops > 0 + assert spec.bandwidth_gb_s > 0 + + +def test_gpu_specs_values_reasonable(): + k100 = GPU_SPECS["K100"] + assert k100.bf16_tflops == 192 + assert k100.bandwidth_gb_s == 700 diff --git a/metainfer/tasks/sglang_trace_analyze/tests/test_plugin.py b/metainfer/tasks/sglang_trace_analyze/tests/test_plugin.py new file mode 100644 index 00000000..357b59d8 --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/tests/test_plugin.py @@ -0,0 +1,18 @@ +"""Validate plugin registration and import sanity.""" + +from metainfer.server.registry import all_plugins +from metainfer.orchestrator.tasks import all_tasks + + +def test_all_plugins_includes_sglang_trace_analyze(): + types = [p.type for p in all_plugins()] + assert "sglang_trace_analyze" in types, ( + f"sglang_trace_analyze not found in registered plugins: {types}" + ) + + +def test_all_tasks_includes_sglang_trace_analyze(): + task_types = [p.task_type for p in all_tasks()] + assert "sglang_trace_analyze" in task_types, ( + f"sglang_trace_analyze not found in registered tasks: {task_types}" + ) diff --git a/metainfer/tasks/sglang_trace_analyze/tests/test_server_readers.py b/metainfer/tasks/sglang_trace_analyze/tests/test_server_readers.py new file mode 100644 index 00000000..385fe8f7 --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/tests/test_server_readers.py @@ -0,0 +1,75 @@ +"""Server state reader tests.""" + +import json +import tempfile +from pathlib import Path + +from ..server._state_readers import ( + read_batch_detail, + read_hints, + read_mapping, + read_summary, +) + + +def test_read_summary(): + with tempfile.TemporaryDirectory() as td: + state_dir = Path(td) + analysis = state_dir / "analysis" + analysis.mkdir() + (analysis / "summary.json").write_text( + json.dumps({"model": "test", "batches": []}) + ) + result = read_summary(state_dir) + assert result is not None + assert result["model"] == "test" + + +def test_read_summary_missing(): + with tempfile.TemporaryDirectory() as td: + assert read_summary(Path(td)) is None + + +def test_read_mapping(): + with tempfile.TemporaryDirectory() as td: + state_dir = Path(td) + analysis = state_dir / "analysis" + analysis.mkdir() + (analysis / "mapping.json").write_text( + json.dumps({"entries": [{"kernel_name": "test"}]}) + ) + result = read_mapping(state_dir) + assert len(result["entries"]) == 1 + + +def test_read_hints(): + with tempfile.TemporaryDirectory() as td: + state_dir = Path(td) + analysis = state_dir / "analysis" + analysis.mkdir() + (analysis / "hints.json").write_text( + json.dumps({"bottleneck": {"kernel_or_pattern": "triton_gemm"}}) + ) + result = read_hints(state_dir) + assert result["bottleneck"]["kernel_or_pattern"] == "triton_gemm" + + +def test_read_batch_detail(): + with tempfile.TemporaryDirectory() as td: + state_dir = Path(td) + batch_dir = state_dir / "analysis" / "batches" / "bs_8" / "decode" + batch_dir.mkdir(parents=True) + (batch_dir / "kernel_table.json").write_text(json.dumps({"kernels": []})) + (batch_dir / "overlap.json").write_text(json.dumps({"gaps": []})) + (batch_dir / "fuse.json").write_text(json.dumps({"matches": []})) + + result = read_batch_detail(state_dir, 8, "decode") + assert result is not None + assert result["kernel_table"]["kernels"] == [] + assert result["overlap"]["gaps"] == [] + + +def test_read_batch_detail_missing(): + with tempfile.TemporaryDirectory() as td: + result = read_batch_detail(Path(td), 8, "decode") + assert result is None diff --git a/metainfer/tasks/sglang_trace_analyze/tests/test_structure_mapper.py b/metainfer/tasks/sglang_trace_analyze/tests/test_structure_mapper.py new file mode 100644 index 00000000..42ec91c5 --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/tests/test_structure_mapper.py @@ -0,0 +1,63 @@ +"""Structure mapper tests.""" + +from ..orchestrator.structure_mapper import ( + _infer_layer, + _infer_op_type, + build_mapping, +) + + +def test_infer_layer_from_call_stack(): + stack = " File \"sglang/srt/layers/attn/triton_ops.py\", line 45\n File \"model.py\"" + layer = _infer_layer(stack, "triton_attn_kernel", {}) + assert "attn" in layer.lower() if layer else True # matched sglang path + + +def test_infer_layer_model_layers_pattern(): + stack = "model.layers.5.self_attn.qkv_proj" + layer = _infer_layer(stack, "triton_gemm", {}) + assert layer == "layer_5" + + +def test_infer_op_type_attention(): + assert _infer_op_type("flash_attn_fwd", "") == "Attention" + assert _infer_op_type("flash_fwd_splitkv_mla", "") == "Attention" + + +def test_infer_op_type_gemm(): + assert _infer_op_type("triton_gemm_kernel", "") == "GEMM" + assert _infer_op_type("w8a8_bf16_matmul", "") == "GEMM" + + +def test_infer_op_type_moe(): + assert _infer_op_type("fused_moe_kernel", "") == "MoE" + + +def test_infer_op_type_norm(): + assert _infer_op_type("rms_norm_kernel", "") == "Norm" + + +def test_infer_op_type_nccl(): + assert _infer_op_type("ncclAllReduce", "") == "NCCL" + + +def test_build_mapping_empty(): + entries = build_mapping([], {"num_hidden_layers": 32}) + assert entries == [] + + +def test_build_mapping_with_call_stack(): + kernels = [ + { + "kernel_name": "triton_gemm", + "total_dur_us": 1000, + "count": 10, + "call_stack": "model.layers.3.self_attn.q_proj", + } + ] + entries = build_mapping(kernels, {"num_hidden_layers": 32}) + assert len(entries) == 1 + assert entries[0]["kernel_name"] == "triton_gemm" + assert entries[0]["model_layer"] == "layer_3" + assert entries[0]["op_type"] == "GEMM" + assert entries[0]["confidence"] == "high" diff --git a/metainfer/tasks/sglang_trace_analyze/tests/test_trace_parser.py b/metainfer/tasks/sglang_trace_analyze/tests/test_trace_parser.py new file mode 100644 index 00000000..c02b1afc --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/tests/test_trace_parser.py @@ -0,0 +1,66 @@ +"""Trace parser tests with synthetic trace fixtures.""" + +import json +from ..orchestrator.trace_parser import aggregate_kernels + + +def _synthetic_trace(kernels): + """Build a minimal Chrome trace JSON.""" + events = [] + for name, dur, extra in kernels: + evt = {"cat": "kernel", "name": name, "ph": "X", "dur": dur, "ts": 0} + if extra: + evt.setdefault("args", {}).update(extra) + events.append(evt) + return {"traceEvents": events} + + +def test_aggregate_empty_trace(): + result = aggregate_kernels(_synthetic_trace([])) + assert result == [] + + +def test_aggregate_single_kernel(): + trace = _synthetic_trace([("triton_gemm", 1000, {})]) + result = aggregate_kernels(trace) + assert len(result) == 1 + assert result[0]["kernel_name"] == "triton_gemm" + assert result[0]["total_dur_us"] == 1000 + assert result[0]["count"] == 1 + + +def test_aggregate_multiple_same_kernel(): + trace = _synthetic_trace([ + ("triton_gemm", 500, {}), + ("triton_gemm", 700, {}), + ("flash_attn", 300, {}), + ]) + result = aggregate_kernels(trace) + assert len(result) == 2 + # triton_gemm aggregates: 500 + 700 = 1200 + assert result[0]["kernel_name"] == "triton_gemm" + assert result[0]["total_dur_us"] == 1200 + assert result[0]["count"] == 2 + # flash_attn is second + assert result[1]["kernel_name"] == "flash_attn" + assert result[1]["total_dur_us"] == 300 + + +def test_aggregate_ignores_non_kernel(): + trace = _synthetic_trace([ + ("triton_gemm", 500, {}), + ("cpu_op", 200, {}), # different cat + ]) + # Make the second event non-kernel + trace["traceEvents"][1]["cat"] = "cpu_op" + result = aggregate_kernels(trace) + assert len(result) == 1 + assert result[0]["kernel_name"] == "triton_gemm" + + +def test_aggregate_includes_call_stack(): + trace = _synthetic_trace([ + ("triton_gemm", 500, {"call stack": "model.layers.5.self_attn"}), + ]) + result = aggregate_kernels(trace, include_call_stack=True) + assert result[0]["call_stack"] == "model.layers.5.self_attn" From 1cacf2192fd1cc1da6368cdd9c087385dda6d995 Mon Sep 17 00:00:00 2001 From: flyingdown Date: Wed, 5 Aug 2026 11:24:51 +0800 Subject: [PATCH 02/15] fix(sglang-trace-analyze): wire analysis modules and complete kernel_table schema - Wire structure_mapper and flops_calculator into pipeline ANALYZE phase - Add all 17 design fields to kernel_table.json (model_layer, tflops_actual, mfu, bound, bandwidth_gb_s, input_dims, confidence) - Fix classifier priority for HIP/CK kernel names (CK-GEMM, CustomAllReduce, MLA, MoE, ElementWise) - Add CPU-op-based model layer inference fallback (no call stacks in trace) - Extract CK GEMM tile dimensions (MTxx) for FLOPs estimation - Add 3-tab frontend: Summary Overview / Batch Detail / Optimization Hints - Register detail_view_module + extra_stylesheets in WebPlugin - Fix trace_parser to accept str paths (not just Path) End-to-end verified: DeepSeek V4 INT8 TP8 BS=8 decode analysis on K100. Co-Authored-By: deepseek-v4-pro[1m] --- .../orchestrator/flops_calculator.py | 30 ++- .../orchestrator/pipeline.py | 170 ++++++++++-- .../orchestrator/structure_mapper.py | 148 ++++++---- .../orchestrator/trace_parser.py | 11 +- .../sglang_trace_analyze/server/plugin.py | 3 + .../sglang_trace_analyze/static/sa-detail.js | 254 ++++++++++++++++++ .../tasks/sglang_trace_analyze/static/sa.css | 65 +++++ 7 files changed, 597 insertions(+), 84 deletions(-) create mode 100644 metainfer/tasks/sglang_trace_analyze/static/sa-detail.js create mode 100644 metainfer/tasks/sglang_trace_analyze/static/sa.css diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/flops_calculator.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/flops_calculator.py index e4192a41..185392de 100644 --- a/metainfer/tasks/sglang_trace_analyze/orchestrator/flops_calculator.py +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/flops_calculator.py @@ -14,6 +14,18 @@ from .gpu_specs import GpuSpec +def extract_ck_tile_dims(kernel_name: str) -> tuple | None: + """Extract (M, N, K) tile dimensions from a CK GEMM kernel name. + + Example: ``Cijk_Alik_Bljk_SB_MT64x128x16_...`` → (64, 128, 16) + """ + import re + m = re.search(r"MT(\d+)x(\d+)x(\d+)", kernel_name) + if m: + return int(m.group(1)), int(m.group(2)), int(m.group(3)) + return None + + def calculate_mfu( kernels: List[Dict[str, Any]], gpu_spec: GpuSpec, @@ -45,29 +57,37 @@ def calculate_mfu( dur_per_invocation_s = dur_s / count if count else dur_s dims = k.get("input_dims", []) op_type = k.get("op_type", "Other") + kernel_name = k.get("kernel_name", "") flops = _estimate_flops(op_type, dims, batch_size) bytes_moved = _estimate_bytes(op_type, dims, batch_size) + # For CK GEMM kernels without input dims, estimate from tile name + if flops == 0 and op_type == "GEMM": + tile = extract_ck_tile_dims(kernel_name) + if tile: + M, N, K_tile = tile + flops = 2 * M * N * K_tile * count + bytes_moved = (M * K_tile + K_tile * N + M * N) * 2 * count + tflops_actual = (flops / dur_s / 1e12) if dur_s > 0 else 0 bandwidth_gb_s = (bytes_moved / dur_s / 1e9) if dur_s > 0 else 0 mfu = (tflops_actual / theoretical_tflops * 100) if theoretical_tflops > 0 else 0 # Compute-bound vs memory-bound heuristic ops_per_byte = flops / bytes_moved if bytes_moved > 0 else float("inf") - # "Roofline" crossover point = peak_flops / peak_bw ops/byte if theoretical_bw > 0: crossover = theoretical_tflops * 1e12 / (theoretical_bw * 1e9) else: crossover = float("inf") bound = "compute" if ops_per_byte > crossover else "memory" - k["tflops_actual"] = round(tflops_actual, 3) + k["tflops_actual"] = round(tflops_actual, 3) if tflops_actual > 0 else None k["tflops_theoretical"] = theoretical_tflops - k["bandwidth_gb_s"] = round(bandwidth_gb_s, 1) + k["bandwidth_gb_s"] = round(bandwidth_gb_s, 1) if bandwidth_gb_s > 0 else None k["bandwidth_theoretical"] = theoretical_bw - k["mfu"] = round(mfu, 1) - k["bound"] = bound + k["mfu"] = round(mfu, 1) if tflops_actual > 0 else None + k["bound"] = bound if (tflops_actual and tflops_actual > 0) else "unknown" k["flops_per_invocation"] = int(flops) return kernels diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/pipeline.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/pipeline.py index 930005da..88906d37 100644 --- a/metainfer/tasks/sglang_trace_analyze/orchestrator/pipeline.py +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/pipeline.py @@ -15,6 +15,7 @@ import re import subprocess import time +from collections import defaultdict from pathlib import Path from typing import Any, Dict, List, Optional @@ -102,6 +103,13 @@ def run(self) -> None: self.store.update_run(current_phase=phase) self.store.append_timeline("phase_enter", {"phase": phase}) + # Skip phases whose outputs already exist (resume / re-run) + if self._phase_is_done(phase): + print(f"[pipeline] phase {phase} output exists, skipping") + self.store.append_timeline("phase_skip", {"phase": phase, "reason": "output exists"}) + phase = next_phase(phase) + continue + method = getattr(self, f"_run_{phase}", None) if method is None: print(f"[pipeline] unknown phase {phase!r}, stopping") @@ -468,6 +476,33 @@ def _run_summarize(self) -> bool: self.store.write_iteration(n, rec.to_dict()) return True + # ================================================================== # + # Phase skip detection (resume / re-run) + # ================================================================== # + + def _phase_is_done(self, phase: str) -> bool: + """Return True if the phase's expected outputs already exist.""" + if phase == "mapping": + return (self._analysis_dir / "mapping.json").exists() + if phase == "benchmark": + # Check that at least one batch_size trace dir exists + for bs in self.batch_sizes: + trace_dir = self.workspace_dir / "traces" / f"bs_{bs}" + if trace_dir.exists(): + return True + return False + if phase == "analyze": + for bs in self.batch_sizes: + for stage in self.stages: + if not (self._analysis_dir / "batches" / f"bs_{bs}" / stage / "kernel_table.json").exists(): + return False + return True + if phase == "hints": + return (self._analysis_dir / "hints.json").exists() + if phase == "summarize": + return (self._analysis_dir / "summary.json").exists() + return False + # ================================================================== # # Helpers # ================================================================== # @@ -517,34 +552,87 @@ def _analyze_one( bs: int, stage: str, ) -> Dict[str, Any]: - """Analyze a single trace file and return kernel_table, overlap, - and fuse results. + """Analyze a single trace file — uses trace_parser, structure_mapper, + flops_calculator, overlap_detector, fuse_matcher.""" + from .trace_parser import parse_trace, aggregate_kernels + from .structure_mapper import _map_one as map_one + from .flops_calculator import calculate_mfu + from .overlap_detector import build_overlap_report + from .fuse_matcher import build_fuse_report + + print(f"[pipeline] loading trace: {trace_path}") + trace_data = parse_trace(str(trace_path)) + + # Aggregate kernels + kernels = aggregate_kernels(trace_data) + total_dur = sum(k["total_dur_us"] for k in kernels) / 1e6 + + # Build CPU op correlation + events = trace_data.get("traceEvents", []) + cpu_ops_by_corr = defaultdict(lambda: []) + kernel_by_corr = defaultdict(lambda: []) + for e in events: + cat = e.get("cat", "") + corr = (e.get("args") or {}).get( + "External id" if cat == "cpu_op" else "correlation" + ) + if cat == "cpu_op" and corr: + cpu_ops_by_corr[corr].append(e.get("name", "")) + elif cat == "kernel" and corr: + kernel_by_corr[corr].append(e.get("name", "")) + + # Map each kernel using structure_mapper + cpu_ops + result_kernels = [] + for k in kernels: + name = k["kernel_name"] + # Collect correlated CPU ops + cpu_ops = set() + for corr, gpu_names in kernel_by_corr.items(): + if name in gpu_names: + for cn in cpu_ops_by_corr.get(corr, []): + cpu_ops.add(cn) + + # Use structure_mapper for op_type/category/layer + mapped = map_one(name, k.get("call_stack", ""), {}, list(cpu_ops)) + pct = k["total_dur_us"] / (total_dur * 1e6) * 100 + + entry = { + "rank": len(result_kernels) + 1, + "kernel_name": name, + "category": mapped["category"], + "op_type": mapped["op_type"], + "model_layer": mapped["model_layer"], + "confidence": mapped["confidence"], + "total_dur_us": k["total_dur_us"], + "time_pct": round(pct, 2), + "count": k["count"], + "avg_dur_us": round(k["total_dur_us"] / k["count"], 2) if k["count"] else 0, + "input_dims": k.get("input_dims", []), + "tflops_theoretical": self.gpu_spec.bf16_tflops, + "bandwidth_theoretical": self.gpu_spec.bandwidth_gb_s, + } + result_kernels.append(entry) + + # Calculate TFLOPS/MFU/bound using flops_calculator + result_kernels = calculate_mfu( + result_kernels, self.gpu_spec, batch_size=bs, dtype="bf16" + ) - Placeholder — real logic in trace_parser / flops_calculator / - overlap_detector / fuse_matcher. - """ - # TODO: implement real analysis pipeline - return { - "kernel_table": { - "model": self.model_path, - "gpu": self.gpu_spec.label, - "batch_size": bs, - "stage": stage, - "kernels": [], - }, - "overlap": { - "batch_size": bs, - "stage": stage, - "gaps": [], - "summary": {"total_gap_us": 0, "total_gap_pct": 0, "cuda_graph_effective": True}, - }, - "fuse": { - "batch_size": bs, - "stage": stage, - "matches": [], - }, + kernel_table = { + "model": self.model_path, + "gpu": self.gpu_spec.label, + "batch_size": bs, + "stage": stage, + "total_gpu_time_s": round(total_dur, 2), + "unique_kernels": len(result_kernels), + "kernels": result_kernels, } + overlap = build_overlap_report(trace_data, bs, stage) + fuse = build_fuse_report(result_kernels, bs, stage) + + return {"kernel_table": kernel_table, "overlap": overlap, "fuse": fuse} + def _llm_generate_hints( self, kernel_summaries: list, @@ -562,3 +650,35 @@ def _llm_generate_hints( "status": "skipped", "reason": "LLM hints not yet wired", } + + +# ------------------------------------------------------------------ # +# Module-level kernel classifier +# ------------------------------------------------------------------ # + +def _classify_kernel(name: str) -> tuple: + """Classify a GPU kernel name into (op_type, category).""" + n = name.lower() + if n.startswith("cijk_"): + return ("GEMM", "CK-GEMM") + if "flash_fwd" in n or "flash_attn" in n: + return ("Attention", "MLA") + if "fused_moe" in n: + return ("MoE", "MoE") + if "nccl" in n: + return ("NCCL", "NCCL-AllGather") + if "allreduce" in n: + return ("NCCL", "NCCL-AllReduce") + if "reduce_kernel" in n: + return ("Reduce", "CustomAllReduce") + if "rms_norm" in n or "rmsnorm" in n: + return ("Norm", "RMSNorm") + if "elementwise" in n: + return ("ElementWise", "ElementWise") + if "gather" in n or "topk" in n: + return ("Memory", "Gather") + if "copy" in n or "memcpy" in n: + return ("Memory", "Copy") + if "vectorized" in n: + return ("ElementWise", "Vectorized") + return ("Other", "Other") diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/structure_mapper.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/structure_mapper.py index 3d5c5dcb..e3f5215b 100644 --- a/metainfer/tasks/sglang_trace_analyze/orchestrator/structure_mapper.py +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/structure_mapper.py @@ -46,14 +46,23 @@ def _map_one( kernel_name: str, call_stack: str, config: Dict[str, Any], + cpu_ops: list | None = None, ) -> Dict[str, Any]: - """Map a single kernel to a model layer by inspecting its call stack.""" - layer = _infer_layer(call_stack, kernel_name, config) - op_type = _infer_op_type(kernel_name, call_stack) + """Map a single kernel to a model layer by inspecting its call stack + and correlated CPU ops.""" + layer = _infer_layer(call_stack, kernel_name, config, cpu_ops) + op_type = _infer_op_type(kernel_name, call_stack, cpu_ops) confidence = "high" if not call_stack: - confidence = "low" + # Without call stacks, we use kernel name + CPU op correlation + has_cpu_hint = bool(cpu_ops) + if has_cpu_hint and _is_ck_gemm(kernel_name): + confidence = "medium" # CK GEMM is unambiguous even without stack + elif has_cpu_hint: + confidence = "medium" + else: + confidence = "low" elif layer is None: confidence = "medium" @@ -67,66 +76,105 @@ def _map_one( } +def _is_ck_gemm(name: str) -> bool: + """CK (composable_kernel) GEMM kernels have Cijk_ prefix.""" + return name.lower().startswith("cijk_") + + def _infer_layer( call_stack: str, kernel_name: str, config: Dict[str, Any], + cpu_ops: list | None = None, ) -> Optional[str]: - """Extract layer information from the call stack. + """Extract layer information from the call stack and kernel name.""" + name_lower = kernel_name.lower() + cpu_lower = " ".join(cpu_ops or []).lower() + + if call_stack: + import re + lines = call_stack.strip().split("\n") + layer_pat = re.compile(r"model\.layers\.(\d+)") + sglang_layer_pat = re.compile( + r"sglang/srt/layers/(attn|moe|mla|linear|norm|embed|sampler|router)" + ) + for line in lines: + m = layer_pat.search(line) + if m: + return f"layer_{m.group(1)}" + m = sglang_layer_pat.search(line) + if m: + return f"layers/{m.group(1)}" + + # Fallback (no call stack): kernel name + CPU op heuristics + if "flash_fwd" in name_lower or "flash_attn" in name_lower: + return "all_layers/attention" + if "fused_moe" in name_lower or "moe" in cpu_lower: + return "moe_layers/experts" + if name_lower.startswith("cijk_"): + return "all_layers/linear" + if "rms_norm" in cpu_lower or "rmsnorm" in name_lower: + return "all_layers/norm" + if "reduce_kernel" in name_lower: + return "all_layers/allreduce" + if "allgather" in cpu_lower or "nccl" in name_lower: + return "all_layers/communication" + if "elementwise" in name_lower or "vectorized" in name_lower: + return "all_layers/elementwise" + + return None - Looks for patterns like: - - ``sglang/srt/layers/...`` - - ``layer_forward`` - - ``model.py``, ``decoder.py``, ``encoder.py`` - - Module names like ``model.layers.5.self_attn`` - Returns ``None`` if no layer info can be inferred. +def _infer_op_type(kernel_name: str, call_stack: str, cpu_ops: list | None = None) -> str: + """Infer the op type from kernel name, call stack, and correlated CPU ops. + + Priority: kernel name patterns > CPU op hints > name substring heuristics. """ - if not call_stack: - return None - - # Heuristic: look for sglang/srt/layers or model.layers.N patterns - lines = call_stack.strip().split("\n") - - # Pattern 1: model.layers.N in the call stack - import re - layer_pat = re.compile(r"model\.layers\.(\d+)") - # Pattern 2: sglang source files under layers/ - sglang_layer_pat = re.compile( - r"sglang/srt/layers/(attn|moe|mla|linear|norm|embed|sampler|router)" - ) - - for line in lines: - m = layer_pat.search(line) - if m: - return f"layer_{m.group(1)}" - m = sglang_layer_pat.search(line) - if m: - return f"layers/{m.group(1)}" - - # Fallback: use kernel name heuristics - if "attn" in kernel_name.lower() or "attention" in kernel_name.lower(): - return "attention (unknown layer)" - if "moe" in kernel_name.lower(): - return "moe (unknown layer)" - if "gemm" in kernel_name.lower() or "linear" in kernel_name.lower(): - return "linear (unknown layer)" + name_lower = kernel_name.lower() + cpu_lower = " ".join(cpu_ops or []).lower() - return None + # ── Strong kernel name patterns (highest priority) ── + + # CK GEMM kernels (HIP/ROCm composable_kernel) + if name_lower.startswith("cijk_"): + return "GEMM" + # GPU kernel name patterns — unambiguous from the kernel name itself + if "nccl" in name_lower: + return "NCCL" + if any(k in name_lower for k in ("flash_fwd", "flash_attn")): + return "Attention" + if "fused_moe" in name_lower: + return "MoE" -def _infer_op_type(kernel_name: str, call_stack: str) -> str: - """Infer the op type from kernel name and call stack.""" - name_lower = kernel_name.lower() - if any(k in name_lower for k in ("attn", "attention", "flash_fwd", "flash_attn")): + # ── Kernel name substring heuristics (medium priority) ── + if "reduce_kernel" in name_lower: + return "Reduce" + if "elementwise" in name_lower: + return "ElementWise" + if "vectorized" in name_lower: + return "ElementWise" + if "gather" in name_lower: + return "Indexing" + + # ── CPU op hints for torch-compiled/fused kernels ── + if "all_reduce" in cpu_lower: + return "Reduce" # CustomAllReduce, not NCCL + if "allgather" in cpu_lower: + return "NCCL" + if "rms_norm" in cpu_lower: + return "Norm" + + # ── Remaining kernel name patterns (lower priority) ── + if any(k in name_lower for k in ("attn", "attention")): return "Attention" - if any(k in name_lower for k in ("moe", "fused_moe")): + if any(k in name_lower for k in ("moe",)): return "MoE" if any(k in name_lower for k in ("gemm", "linear", "matmul", "w8a8", "fp8")): return "GEMM" - if any(k in name_lower for k in ("rms", "norm", "layernorm", "layer_norm")): + if any(k in name_lower for k in ("rmsnorm", "rms_norm", "layernorm")): return "Norm" - if any(k in name_lower for k in ("nccl", "allreduce", "allgather", "broadcast")): + if any(k in name_lower for k in ("allreduce", "allgather", "broadcast")): return "NCCL" if any(k in name_lower for k in ("hadamard", "rotate", "rope")): return "Transform" @@ -134,7 +182,7 @@ def _infer_op_type(kernel_name: str, call_stack: str) -> str: return "Memory" if any(k in name_lower for k in ("silu", "gelu", "swiglu", "activation", "act_and_mul")): return "Activation" - if any(k in name_lower for k in ("topk", "top_k", "index", "gather", "scatter", "sort")): + if any(k in name_lower for k in ("topk", "top_k", "gather", "scatter", "sort")): return "Indexing" if any(k in name_lower for k in ("quant", "dequant", "fp8_scale")): return "Quantization" @@ -154,5 +202,7 @@ def _op_type_to_category(op_type: str) -> str: "Activation": "Activation", "Indexing": "Indexing", "Quantization": "Quantization", + "Reduce": "Reduce", + "ElementWise": "ElementWise", } return mapping.get(op_type, "Other") diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/trace_parser.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/trace_parser.py index 83bc8a58..7abf9dc0 100644 --- a/metainfer/tasks/sglang_trace_analyze/orchestrator/trace_parser.py +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/trace_parser.py @@ -18,14 +18,15 @@ from typing import Any, Dict, List, Optional -def _open_trace(trace_path: Path): +def _open_trace(trace_path): """Open a trace file — transparently handles .gz compression.""" - if trace_path.suffix == ".gz": - return gzip.open(trace_path, "rt", encoding="utf-8") - return open(trace_path, "r", encoding="utf-8") + tp = Path(trace_path) + if tp.suffix == ".gz": + return gzip.open(tp, "rt", encoding="utf-8") + return open(tp, "r", encoding="utf-8") -def parse_trace(trace_path: Path) -> Dict[str, Any]: +def parse_trace(trace_path) -> Dict[str, Any]: """Load a Chrome trace JSON and return the top-level document. Returns: diff --git a/metainfer/tasks/sglang_trace_analyze/server/plugin.py b/metainfer/tasks/sglang_trace_analyze/server/plugin.py index 6fdaa277..ea1816f8 100644 --- a/metainfer/tasks/sglang_trace_analyze/server/plugin.py +++ b/metainfer/tasks/sglang_trace_analyze/server/plugin.py @@ -24,8 +24,11 @@ "optimization hints." ), build_router=build_router, + detail_view_module="app/sa-detail", + detail_view_export="default", frontend_dir=_FRONTEND_DIR, importmap_entries=_IMPORTMAP_ENTRIES, + extra_stylesheets=["sa.css"], ) register(plugin) diff --git a/metainfer/tasks/sglang_trace_analyze/static/sa-detail.js b/metainfer/tasks/sglang_trace_analyze/static/sa-detail.js new file mode 100644 index 00000000..710bf03d --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/static/sa-detail.js @@ -0,0 +1,254 @@ +/** SGLang Trace Analyze — task detail view. + * + * Three tabs: Summary Overview | Batch Detail | Optimization Hints + */ +import { html } from "htm/preact"; +import { useCallback, useEffect, useState } from "preact/hooks"; + +const API = (taskId) => `/api/sglang_trace_analyze/${taskId}`; + +export default function SADetail({ taskId }) { + const [summary, setSummary] = useState(null); + const [hints, setHints] = useState(null); + const [detail, setDetail] = useState(null); + const [activeTab, setActiveTab] = useState("summary"); + const [activeBatch, setActiveBatch] = useState(null); + const [activeStage, setActiveStage] = useState("decode"); + const [loading, setLoading] = useState(true); + const [error, setError] = useState(null); + + useEffect(() => { + Promise.all([ + fetch(`${API(taskId)}/summary`).then((r) => r.json()), + fetch(`${API(taskId)}/hints`).then((r) => r.json()), + ]) + .then(([s, h]) => { setSummary(s); setHints(h); setLoading(false); }) + .catch((e) => { setError(e.message); setLoading(false); }); + }, [taskId]); + + useEffect(() => { + if (!activeBatch) return; + fetch(`${API(taskId)}/batch/${activeBatch}/${activeStage}`) + .then((r) => r.json()) + .then((d) => setDetail(d)) + .catch(() => setDetail(null)); + }, [taskId, activeBatch, activeStage]); + + if (loading) return html`
Loading analysis…
`; + if (error) return html`
Error: ${error}
`; + if (!summary || !summary.batches || summary.batches.length === 0) { + return html`
No analysis data available yet.
`; + } + + const batchList = summary.batches || []; + if (!activeBatch && batchList.length > 0) { + setActiveBatch(batchList[0].batch_size); + } + + return html` +
+
+

SGLang Trace Analysis

+ Model: ${summary.model || "?"} | GPU: ${summary.gpu || "?"} +
+ +
+ + + +
+ + ${activeTab === "summary" && html`<${SummaryPage} summary=${summary} />`} + ${activeTab === "batch" && html` +
+ ${batchList.map((b) => html` + + `)} +
+ ${detail ? html` + <${KernelTable} kt=${detail.kernel_table} batch=${activeBatch} stage=${activeStage} /> + <${CategoryChart} kt=${detail.kernel_table} /> + <${OverlapPanel} ov=${detail.overlap} fu=${detail.fuse} /> + ` : html`
Loading batch detail…
`} + `} + ${activeTab === "hints" && html`<${HintsPage} hints=${hints} />`} +
+ `; +} + +/* ── Summary Overview ── */ + +function SummaryPage({ summary }) { + const batches = summary.batches || []; + return html` +
+

Batch Summary

+ + + + + + ${batches.map((b) => html` + + + + + + + + + `)} + +
BatchStageTop Kernel%KernelsMFU Avg
${b.batch_size}${b.stage}${(b.top_kernel || "").slice(0, 70)}${(b.top_kernel_pct || 0).toFixed(1)}%${b.kernel_count}${b.mfu_avg != null ? b.mfu_avg.toFixed(1) + "%" : "-"}
+
+ `; +} + +/* ── Kernel Hotspot Table ── */ + +function KernelTable({ kt, batch, stage }) { + if (!kt) return null; + const kernels = kt.kernels || []; + const totalTime = kt.total_gpu_time_s || 1; + + return html` +
+

Kernel Hotspots — BS=${batch} ${stage} (${kernels.length} unique, ${totalTime.toFixed(1)}s GPU)

+
+ + + + + + ${kernels.slice(0, 25).map((k) => html` + + + + + + + + + + + + + `)} + +
#%CategoryOpLayerMFUBoundCountAvg μsKernel
${k.rank}
${(k.time_pct || 0).toFixed(1)}%
${k.category || "?"}${k.op_type || "?"}${k.model_layer || "-"}${k.mfu != null ? k.mfu.toFixed(1) + "%" : "-"}${k.bound || "-"}${k.count}${(k.avg_dur_us || 0).toFixed(1)}${(k.kernel_name || "").slice(0, 60)}
+
+
+ `; +} + +/* ── Category Breakdown ── */ + +function CategoryChart({ kt }) { + if (!kt) return null; + const kernels = kt.kernels || []; + const cats = {}; + for (const k of kernels) { + cats[k.category] = (cats[k.category] || 0) + (k.time_pct || 0); + } + const sorted = Object.entries(cats).sort((a, b) => b[1] - a[1]); + const colors = ["#c0392b","#d35400","#e67e22","#27ae60","#2980b9","#8e44ad","#16a085","#7f8c8d","#2c3e50","#e91e63"]; + + return html` +
+

Category Breakdown

+
+ ${sorted.map(([cat, pct], i) => html` +
+ ${cat} +
+
+
+ ${pct.toFixed(1)}% +
+ `)} +
+
+ `; +} + +/* ── Overlap + Fuse ── */ + +function OverlapPanel({ ov, fu }) { + if (!ov) return null; + const sum = ov.summary || {}; + const gaps = ov.gaps || []; + const high = gaps.filter((g) => g.severity === "high").length; + const medium = gaps.filter((g) => g.severity === "medium").length; + const fuseMatches = (fu && fu.matches) || []; + + return html` +
+

Overlap & Fuse

+

+ ${gaps.length} GPU idle gaps (${high} high, ${medium} medium) + — total: ${(sum.total_gap_us / 1000).toFixed(1)}ms + (${sum.total_gap_pct || 0}% of GPU time) +

+

CUDA Graph effective: ${sum.cuda_graph_effective ? "YES" : "NO"}

+ ${fuseMatches.length > 0 && html` +

Fuse Pattern Matches

+ ${fuseMatches.map((m) => html` +
+ ${m.pattern} (${m.confidence}) +

${m.suggestion}

+
+ `)} + `} +
+ `; +} + +/* ── Hints Page ── */ + +function HintsPage({ hints }) { + if (!hints) return html`

No hints generated yet.

`; + + const b = hints.bottleneck || {}; + const suggestions = hints.suggestions || []; + const surprises = hints.surprises || []; + + if (hints.status === "skipped") { + return html`

Optimization Hints

Hints generation skipped (${hints.reason || "not wired"}).

`; + } + + return html` +
+

Optimization Hints

+
+

Biggest Bottleneck

+

${b.kernel_or_pattern || "?"} — ${b.reason || ""} (impact: ${b.impact_pct || 0}%)

+
+ ${suggestions.length > 0 && html` +
+

Suggestions

+ ${suggestions.map((s) => html` +
+ ${s.title} ${s.difficulty} +

${s.what_to_change}

+

Why: ${s.why} | Est. saving: ${s.estimated_saving_pct}% | Type: ${s.category}

+
+ `)} +
+ `} + ${surprises.length > 0 && html` +
+

Surprises

+ ${surprises.map((s) => html`

${s}

`)} +
+ `} +
+ `; +} diff --git a/metainfer/tasks/sglang_trace_analyze/static/sa.css b/metainfer/tasks/sglang_trace_analyze/static/sa.css new file mode 100644 index 00000000..ca23dcb2 --- /dev/null +++ b/metainfer/tasks/sglang_trace_analyze/static/sa.css @@ -0,0 +1,65 @@ +.sa-detail { padding: 12px 16px; color: #e0e0e0; font-family: system-ui, sans-serif; } +.sa-header { margin-bottom: 12px; } +.sa-header h2 { margin: 0 0 2px; color: #fff; font-size: 18px; } +.sa-meta { color: #888; font-size: 12px; } +.sa-loading,.sa-error,.sa-empty { padding: 32px; text-align: center; color: #888; } +.sa-error { color: #e74c3c; } + +/* Tabs */ +.sa-tabs { display: flex; gap: 2px; margin-bottom: 12px; border-bottom: 2px solid #333; } +.sa-tab-btn { padding: 6px 16px; border: none; border-radius: 4px 4px 0 0; + background: transparent; color: #999; cursor: pointer; font-size: 13px; } +.sa-tab-btn.active-tab { background: #2a2a2a; color: #4a90d9; font-weight: 600; } +.sa-batch-tabs { display: flex; gap: 6px; margin-bottom: 12px; } +.sa-tab { padding: 4px 12px; border: 1px solid #444; border-radius: 4px; + background: #2a2a2a; color: #ccc; cursor: pointer; font-size: 12px; } +.sa-tab.active { background: #4a90d9; color: #fff; border-color: #4a90d9; } + +/* Panels */ +.sa-panel { background: #1e1e1e; border: 1px solid #333; border-radius: 6px; + padding: 14px; margin-bottom: 12px; } +.sa-panel h3 { margin: 0 0 10px; color: #ddd; font-size: 14px; } +.sa-note { color: #888; font-size: 11px; margin: 4px 0; } + +/* Table */ +.sa-table-wrap { overflow-x: auto; } +.sa-table { width: 100%; border-collapse: collapse; font-size: 11px; } +.sa-table th { text-align: left; padding: 5px 6px; border-bottom: 1px solid #333; + color: #999; font-weight: 600; white-space: nowrap; } +.sa-table td { padding: 3px 6px; border-bottom: 1px solid #2a2a2a; vertical-align: middle; } +.sa-num { text-align: right; font-variant-numeric: tabular-nums; color: #aaa; } +.sa-sm { font-size: 10px; color: #888; } +.sa-kernel-name { max-width: 280px; overflow: hidden; text-overflow: ellipsis; + white-space: nowrap; font-family: monospace; font-size: 10px; color: #bbb; } +.sa-pct { width: 100px; } +.sa-bar-bg { position: relative; background: #2a2a2a; border-radius: 2px; + height: 14px; overflow: hidden; } +.sa-bar { position: absolute; left: 0; top: 0; height: 100%; + background: #4a90d9; border-radius: 2px; opacity: 0.5; } +.sa-bar-bg span { position: relative; z-index: 1; font-size: 10px; + line-height: 14px; padding-left: 3px; color: #ddd; } +.sa-cat { display: inline-block; padding: 1px 5px; border-radius: 2px; + font-size: 10px; background: #333; color: #ccc; } + +/* Category chart */ +.sa-cat-chart { display: flex; flex-direction: column; gap: 5px; } +.sa-cat-row { display: flex; align-items: center; gap: 8px; } +.sa-cat-label { width: 120px; font-size: 11px; color: #ccc; text-align: right; } +.sa-cat-bar-bg { flex: 1; background: #2a2a2a; border-radius: 2px; height: 16px; overflow: hidden; } +.sa-cat-bar { height: 100%; border-radius: 2px; min-width: 2px; } +.sa-cat-pct { width: 50px; font-size: 11px; color: #aaa; font-variant-numeric: tabular-nums; } + +/* Hints */ +.sa-hint-section { margin-top: 10px; } +.sa-hint-section h4 { margin: 0 0 6px; color: #ccc; font-size: 12px; } +.sa-hint-card { background: #252525; border-left: 3px solid #4a90d9; + padding: 8px 10px; margin-bottom: 8px; border-radius: 0 4px 4px 0; } +.sa-difficulty { display: inline-block; padding: 1px 6px; border-radius: 3px; + font-size: 10px; color: #fff; margin-left: 6px; } +.sa-diff-low { background: #27ae60; } +.sa-diff-medium { background: #e67e22; } +.sa-diff-high { background: #c0392b; } + +/* Fuse cards */ +.sa-fuse-card { background: #252525; padding: 6px 8px; margin: 6px 0; border-radius: 4px; } +.sa-fuse-card strong { color: #f1c40f; } From 1f837d07702b9fd8d45d41290c05d391267a5ced Mon Sep 17 00:00:00 2001 From: flyingdown Date: Wed, 5 Aug 2026 13:32:01 +0800 Subject: [PATCH 03/15] fix(sglang-trace-analyze): add DeepSeek V4 parser flags to run_benchmark MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Add --reasoning-parser deepseek-v4 and --tool-call-parser deepseekv4 to match upstream /workspace/sglang/scripts/run_traces.py params. Without these, sglang may use incorrect model config parser on startup. Verified: upstream run_traces.sh also SIGSEGVs with CUDA Graph ON on K100 — this is a sglang fork bug, not a parameter mismatch. Co-Authored-By: deepseek-v4-pro[1m] --- .../tasks/sglang_trace_analyze/orchestrator/run_benchmark.py | 2 ++ 1 file changed, 2 insertions(+) diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/run_benchmark.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/run_benchmark.py index 469cf07e..fb77475e 100644 --- a/metainfer/tasks/sglang_trace_analyze/orchestrator/run_benchmark.py +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/run_benchmark.py @@ -66,6 +66,8 @@ def run_benchmark( "--chunked-prefill-size", "4096", "--kv-cache-dtype", "auto", "--disable-flashinfer-autotune", + "--reasoning-parser", "deepseek-v4", + "--tool-call-parser", "deepseekv4", "--enable-metrics", ] From 4e7d36d09e1cf2653dbb91e17efef930bc3a9195 Mon Sep 17 00:00:00 2001 From: flyingdown Date: Wed, 5 Aug 2026 15:54:49 +0800 Subject: [PATCH 04/15] fix(sglang-trace-analyze): make pipeline resilient to BENCHMARK failure MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit - Real _build_mapping: parse trace with CPU op correlation, call structure_mapper to classify all kernel→layer mappings - Configurable profile steps in bench_config.json (not hardcoded 500/50) - BENCHMARK failure is non-fatal: ANALYZE falls back to mapping traces - ANALYZE auto-discovers traces in sglang timestamp subdirectories - run_benchmark.py uses configurable profile_start_step/profile_steps Verified: full 5-phase pipeline (MAPPING→BENCHMARK→ANALYZE→HINTS→ SUMMARIZE) completes with final_status=success even when formal benchmark SIGSEGVs (sglang K100 fork CUDA Graph bug). Co-Authored-By: deepseek-v4-pro[1m] --- .../orchestrator/pipeline.py | 123 ++++++++++++++---- .../orchestrator/run_benchmark.py | 4 +- 2 files changed, 97 insertions(+), 30 deletions(-) diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/pipeline.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/pipeline.py index 88906d37..cdce8387 100644 --- a/metainfer/tasks/sglang_trace_analyze/orchestrator/pipeline.py +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/pipeline.py @@ -124,9 +124,14 @@ def run(self) -> None: return if not ok: - print(f"[pipeline] phase {phase} returned failure, stopping") - self.store.update_run(finished=True, final_status="failed") - return + # MAPPING failure is fatal (no traces to analyze). + # BENCHMARK failure is non-fatal: ANALYZE can still use + # MAPPING traces (CUDA Graph OFF) with a note. + if phase == "mapping": + print(f"[pipeline] phase {phase} returned failure, stopping") + self.store.update_run(finished=True, final_status="failed") + return + print(f"[pipeline] phase {phase} returned failure, continuing with available data") self.store.append_timeline("phase_exit", {"phase": phase}) phase = next_phase(phase) @@ -182,19 +187,35 @@ def _run_mapping(self) -> bool: return False # 3. Parse trace → build mapping table (rule engine) - decode_trace_dir = trace_dir / "decode" + # sglang puts traces inside a timestamp subdirectory + decode_trace_dir = trace_dir if not decode_trace_dir.exists(): - # try: the wrapper may have used --profile-by-stage naming - candidates = sorted(trace_dir.glob("*.json.gz")) - if not candidates: - rec.fail("no trace files found after mapping benchmark") - self.store.write_iteration(n, rec.to_dict()) - return False - trace_path = candidates[0] # best effort + # try globbing for timestamp subdirs + ts_dirs = sorted(trace_dir.parent.glob( + trace_dir.name + "/*" if trace_dir.name else "*/" + )) if trace_dir.parent.exists() else [] + if not ts_dirs: + # fall back: find any trace files + candidates = list(trace_dir.parent.rglob("*.trace.json.gz")) if trace_dir.parent.exists() else [] + if not candidates: + rec.fail("no trace files found after mapping benchmark") + self.store.write_iteration(n, rec.to_dict()) + return False + trace_path = candidates[0] + else: + decode_trace_dir = ts_dirs[0] + traces = sorted(decode_trace_dir.glob("*DECODE*.trace.json.gz")) + if not traces: + rec.fail("no decode traces in " + str(decode_trace_dir)) + self.store.write_iteration(n, rec.to_dict()) + return False + trace_path = traces[0] else: - traces = sorted(decode_trace_dir.glob("*.trace.json.gz")) + traces = sorted(decode_trace_dir.glob("*DECODE*.trace.json.gz")) + if not traces: + traces = sorted(decode_trace_dir.rglob("*DECODE*.trace.json.gz")) if not traces: - rec.fail("no trace files in mapping/decode/") + rec.fail("no decode traces found in " + str(decode_trace_dir)) self.store.write_iteration(n, rec.to_dict()) return False trace_path = traces[0] @@ -286,11 +307,13 @@ def _run_benchmark(self) -> bool: rec.done(trace_dir=str(trace_dir)) self.store.write_iteration(n, rec.to_dict()) - # We continue even if some batches failed — ANALYZE skips them - return all_ok or any( - (self.workspace_dir / "traces" / f"bs_{bs}" / "decode").exists() + # Continue if any traces exist (mapping or formal) + has_formal = any( + (self.workspace_dir / "traces" / f"bs_{bs}").exists() for bs in self.batch_sizes ) + has_mapping = (self.workspace_dir / "traces" / "mapping").exists() + return all_ok or has_formal or has_mapping # ================================================================== # # Phase: ANALYZE @@ -308,12 +331,22 @@ def _run_analyze(self) -> bool: any_ok = False for bs in self.batch_sizes: for stage in self.stages: + # Try formal traces first, fall back to mapping traces trace_dir = self.workspace_dir / "traces" / f"bs_{bs}" / stage if not trace_dir.exists(): - print(f"[pipeline] skipping bs_{bs}/{stage} — no trace dir") - continue - - traces = sorted(trace_dir.glob("*.trace.json.gz")) + # Fallback: look for mapping trace subdir + map_base = self.workspace_dir / "traces" / "mapping" + if map_base.exists(): + ts_dirs = sorted(map_base.glob("*/")) # timestamp subdirs + if ts_dirs: + trace_dir = ts_dirs[0] + else: + trace_dir = map_base + else: + print(f"[pipeline] skipping bs_{bs}/{stage} — no trace dir") + continue + + traces = sorted(trace_dir.glob("*DECODE*.trace.json.gz")) if not traces: print(f"[pipeline] skipping bs_{bs}/{stage} — no trace files") continue @@ -520,18 +553,52 @@ def _build_bench_config( "tp_size": self.tp_size, "pp_size": self.pp_size, "output_dir": output_dir, + "profile_start_step": 5, + "profile_steps": 5, } def _build_mapping(self, trace_path: Path) -> List[Dict[str, Any]]: - """Parse a torch profiler Chrome trace and extract kernel→layer - mappings from call stacks. + """Parse a trace file and build kernel→model-structure mapping + using trace_parser + structure_mapper with CPU op correlation.""" + from .trace_parser import parse_trace, aggregate_kernels + from .structure_mapper import _map_one as map_one - Placeholder implementation — real logic will live in - ``trace_parser.py`` and ``structure_mapper.py``. - """ - # TODO: implement trace_parser.py - print("[pipeline] _build_mapping: parsing trace (placeholder)") - return [] + print(f"[pipeline] parsing trace for mapping: {trace_path}") + trace_data = parse_trace(str(trace_path)) + + kernels = aggregate_kernels(trace_data) + + # Build CPU op correlation + events = trace_data.get("traceEvents", []) + cpu_ops_by_corr = defaultdict(lambda: []) + kernel_by_corr = defaultdict(lambda: []) + for e in events: + cat = e.get("cat", "") + corr = (e.get("args") or {}).get( + "External id" if cat == "cpu_op" else "correlation" + ) + if cat == "cpu_op" and corr: + cpu_ops_by_corr[corr].append(e.get("name", "")) + elif cat == "kernel" and corr: + kernel_by_corr[corr].append(e.get("name", "")) + + entries = [] + seen = set() + for k in kernels: + name = k["kernel_name"] + if name in seen: + continue + seen.add(name) + cpu_ops = set() + for corr, gpu_names in kernel_by_corr.items(): + if name in gpu_names: + for cn in cpu_ops_by_corr.get(corr, []): + cpu_ops.add(cn) + entry = map_one(name, k.get("call_stack", ""), {}, list(cpu_ops)) + entries.append(entry) + + print(f"[pipeline] mapping built: {len(entries)} unique kernels") + return entries def _llm_mapping_sanity_check( self, entries: List[Dict[str, Any]] diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/run_benchmark.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/run_benchmark.py index fb77475e..52e4edeb 100644 --- a/metainfer/tasks/sglang_trace_analyze/orchestrator/run_benchmark.py +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/run_benchmark.py @@ -57,8 +57,8 @@ def run_benchmark( "--dataset-name", "random-ids", "--fake-prefill", "--profile", - "--profile-start-step", "500", - "--profile-steps", "50", + "--profile-start-step", str(args.get("profile_start_step", 5)), + "--profile-steps", str(args.get("profile_steps", 5)), "--profile-by-stage", "--profile-prefix", profile_prefix, "--profile-output-dir", output_dir, From c772519592af56c32a4b9aed039a1b0f0fe449ad Mon Sep 17 00:00:00 2001 From: flyingdown Date: Wed, 5 Aug 2026 17:18:19 +0800 Subject: [PATCH 05/15] fix(sglang-trace-analyze): correct build_dir_name to reflect CUDA Graph state Always appended "graph" regardless of --disable-cuda-graph. Now uses "nograph" for mapping runs, matching upstream run_traces.py behavior. Co-Authored-By: deepseek-v4-pro[1m] --- .../sglang_trace_analyze/orchestrator/run_benchmark.py | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/run_benchmark.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/run_benchmark.py index 52e4edeb..7b81d833 100644 --- a/metainfer/tasks/sglang_trace_analyze/orchestrator/run_benchmark.py +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/run_benchmark.py @@ -24,10 +24,10 @@ from typing import Dict, Any, List -def build_dir_name(args: Dict[str, Any]) -> str: +def build_dir_name(args: Dict[str, Any], disable_cuda_graph: bool = False) -> str: """Build sglang-style directory name from config.""" parts = [args["version"], f"tp{args['tp_size']}", f"pp{args['pp_size']}"] - parts.append("graph") + parts.append("nograph" if disable_cuda_graph else "graph") return "_".join(parts) @@ -116,14 +116,14 @@ def main(): with open(config_path) as f: cfg = json.load(f) - dir_name = build_dir_name(cfg) - if args.mapping_only: + dir_name = build_dir_name(cfg, disable_cuda_graph=True) bs = cfg.get("mapping_batch_size", 8) ok = run_benchmark(cfg, dir_name, bs, disable_cuda_graph=True) return 0 if ok else 1 if args.formal_only: + dir_name = build_dir_name(cfg, disable_cuda_graph=False) batch_sizes: List[int] = cfg.get("batch_sizes", [1]) if args.single_batch is not None: if args.single_batch in batch_sizes: From 6ba0d95343038c8a329658bee9a7a9e8266a38e6 Mon Sep 17 00:00:00 2001 From: flyingdown Date: Wed, 5 Aug 2026 20:43:05 +0800 Subject: [PATCH 06/15] fix(sglang-trace-analyze): classify cross_device_reduce as Reduce MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit vllm::cross_device_reduce_2stage_pcie is the TP allreduce kernel on K100 HIP. Previously classified as Other (63.8% of GPU time with CUDA Graph ON), now correctly classified as Reduce. Verified: CUDA Graph bug fixed upstream, formal run produces traces. GPU time drops from 7.09s (no graph) to 0.57s (graph ON, 12.4x), throughput 5.63 → 19.71 tok/s (3.5x). Bottleneck shifts from Reduce (71.9%) to GEMM (46.5%). Co-Authored-By: deepseek-v4-pro[1m] --- .../tasks/sglang_trace_analyze/orchestrator/structure_mapper.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/structure_mapper.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/structure_mapper.py index e3f5215b..6eb4af0d 100644 --- a/metainfer/tasks/sglang_trace_analyze/orchestrator/structure_mapper.py +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/structure_mapper.py @@ -148,7 +148,7 @@ def _infer_op_type(kernel_name: str, call_stack: str, cpu_ops: list | None = Non return "MoE" # ── Kernel name substring heuristics (medium priority) ── - if "reduce_kernel" in name_lower: + if "reduce_kernel" in name_lower or "cross_device_reduce" in name_lower: return "Reduce" if "elementwise" in name_lower: return "ElementWise" From a11de71cb4ea4e33e1c1e0c2912db51e7ea9b1a6 Mon Sep 17 00:00:00 2001 From: flyingdown Date: Thu, 6 Aug 2026 00:36:59 +0800 Subject: [PATCH 07/15] fix(sglang-trace-analyze): embed K100 env vars into run_benchmark.py MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Previously required caller to source HIP/K100 environment variables externally. Now _setup_env() applies them at startup, matching /workspace/sglang/scripts/run_traces.sh exactly. Key fixes: - SGLANG_OPT_USE_HIP_INT8_SCALED_MM: true → 0 - Added SGLANG_OPT_USE_LMSLIM_INT8_QUANT=1 - Added SGLANG_OPT_USE_W8A8_MARLIN_GEMM=1 Co-Authored-By: deepseek-v4-pro[1m] --- .../orchestrator/run_benchmark.py | 57 ++++++++++++++++++- 1 file changed, 54 insertions(+), 3 deletions(-) diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/run_benchmark.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/run_benchmark.py index 7b81d833..9f90d079 100644 --- a/metainfer/tasks/sglang_trace_analyze/orchestrator/run_benchmark.py +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/run_benchmark.py @@ -9,8 +9,8 @@ # Formal runs — one or all batch sizes, CUDA Graph ON python run_benchmark.py --config bench_config.json --formal-only [--single-batch N] -The benchmark is a synchronous, blocking call — the caller waits for -all batch sizes to complete. +Environment variables required for K100/HIP are set inside this script +so callers don't need to source them externally. """ from __future__ import annotations @@ -21,7 +21,57 @@ import subprocess import sys from pathlib import Path -from typing import Dict, Any, List +from typing import Any, Dict, List + +# ── K100 / HIP environment — must match /workspace/sglang/scripts/run_traces.sh ── + +_K100_ENV = { + "SGL_CHUNKED_PREFIX_CACHE_THRESHOLD": "0", + "SGLANG_DISAGGREGATION_BOOTSTRAP_TIMEOUT": "1200", + "GLIBC_TUNABLES": "glibc.rtld.optional_static_tls=0x40000", + "SGLANG_SET_CPU_AFFINITY": "1", + "HIP_KERNEL_BATCH_CEILING": "100", + "GPU_MAX_HW_QUEUES": "3", + # "HIP_GRAPH_ACCUMULATE_DISPATCH": "0", # torchprof needs this + "HIP_H2D_DISABLE_COPY_BUFFER": "0", + "HIP_D2H_DISABLE_COPY_BUFFER": "0", + "HIP_H2D_DIRECT_COPY_THRESHOLD": "32768", + "HIP_H2D_HSAAPI_COPY_THRESHOLD": "32768", + "HIP_D2H_DIRECT_COPY_THRESHOLD": "512", + "HIP_D2H_HSAAPI_COPY_THRESHOLD": "512", + "USE_DCU_CUSTOM_ALLREDUCE": "1", + "HIP_KERNEL_EVENT_SYSTENFENCE": "1", + "SGLANG_USE_FP8_W8A8_MOE": "0", + "SGLANG_USE_LIGHTOP": "0", + "SGLANG_ROCM_USE_AITER_MOE": "0", + "SGLANG_OPT_USE_FUSED_HASH_TOPK": "false", + "SGLANG_OPT_SWIGLU_CLAMP_FUSION": "false", + "SGLANG_TOPK_TRANSFORM_512_TORCH": "false", + "SGLANG_OPT_USE_JIT_KERNEL_FUSED_TOPK": "false", + "SGLANG_NSA_FUSE_TOPK": "false", + "SGLANG_JIT_DEEPGEMM_PRECOMPILE": "0", + "SGLANG_APPLY_CONFIG_BACKUP": "none", + "SGLANG_DSV4_MODE": "2604", + "SGLANG_OPT_BF16_FP32_GEMM_ALGO": "torch", + "SGLANG_OPT_USE_HIP_PAGED_MQA_LOGITS": "1", + "SGLANG_OPT_USE_HIP_MHC_PRE": "1", + "SGLANG_OPT_USE_HIP_MHC_POST": "1", + "SGLANG_OPT_USE_HIP_INT8_SCALED_MM": "0", + "SGLANG_OPT_USE_LMSLIM_INT8_QUANT": "1", + "SGLANG_OPT_USE_W8A8_MARLIN_GEMM": "1", +} + +_PYTHONPATH_EXTRA = "/workspace/sglang/sglang-v0.5.15_k100/python" + + +def _setup_env(): + """Apply K100 env vars and PYTHONPATH once per process.""" + for k, v in _K100_ENV.items(): + if k not in os.environ: + os.environ[k] = v + pp = os.environ.get("PYTHONPATH", "") + if _PYTHONPATH_EXTRA not in pp: + os.environ["PYTHONPATH"] = f"{_PYTHONPATH_EXTRA}:{pp}" if pp else _PYTHONPATH_EXTRA def build_dir_name(args: Dict[str, Any], disable_cuda_graph: bool = False) -> str: @@ -95,6 +145,7 @@ def run_benchmark( def main(): + _setup_env() parser = argparse.ArgumentParser( description="Run sglang bench_one_batch_server with torch profiler" ) From 2f9678ff77f3070105f9b0eeed6796c861a597cf Mon Sep 17 00:00:00 2001 From: flyingdown Date: Thu, 6 Aug 2026 00:39:16 +0800 Subject: [PATCH 08/15] feat(sglang-trace-analyze): redesign frontend with dashboard, donut chart, search Replace single-page layout with 3-tab dashboard: - Dashboard: stat cards (GPU time, bottleneck %, MFU, CUDA Graph), CSS donut chart for category breakdown, bottleneck detail card, compute/memory bound visualization, overlap status, top kernels preview - Kernel Table: search bar + category filter, all 11 columns with confidence badges, sortable and filterable - Hints: bottleneck analysis with auto-generated suggestions, fuse pattern matches, AI optimization hints Co-Authored-By: deepseek-v4-pro[1m] --- .../sglang_trace_analyze/static/sa-detail.js | 420 +++++++++++------- .../tasks/sglang_trace_analyze/static/sa.css | 102 ++++- 2 files changed, 349 insertions(+), 173 deletions(-) diff --git a/metainfer/tasks/sglang_trace_analyze/static/sa-detail.js b/metainfer/tasks/sglang_trace_analyze/static/sa-detail.js index 710bf03d..58a7f3f3 100644 --- a/metainfer/tasks/sglang_trace_analyze/static/sa-detail.js +++ b/metainfer/tasks/sglang_trace_analyze/static/sa-detail.js @@ -1,9 +1,10 @@ /** SGLang Trace Analyze — task detail view. * - * Three tabs: Summary Overview | Batch Detail | Optimization Hints + * Designed for GPU inference optimization engineers. + * Three tabs: Dashboard | Batch Detail | Optimization Hints */ import { html } from "htm/preact"; -import { useCallback, useEffect, useState } from "preact/hooks"; +import { useEffect, useState, useMemo } from "preact/hooks"; const API = (taskId) => `/api/sglang_trace_analyze/${taskId}`; @@ -11,7 +12,7 @@ export default function SADetail({ taskId }) { const [summary, setSummary] = useState(null); const [hints, setHints] = useState(null); const [detail, setDetail] = useState(null); - const [activeTab, setActiveTab] = useState("summary"); + const [activeTab, setActiveTab] = useState("dashboard"); const [activeBatch, setActiveBatch] = useState(null); const [activeStage, setActiveStage] = useState("decode"); const [loading, setLoading] = useState(true); @@ -41,214 +42,325 @@ export default function SADetail({ taskId }) { } const batchList = summary.batches || []; - if (!activeBatch && batchList.length > 0) { - setActiveBatch(batchList[0].batch_size); - } + if (!activeBatch && batchList.length > 0) setActiveBatch(batchList[0].batch_size); return html`
-

SGLang Trace Analysis

- Model: ${summary.model || "?"} | GPU: ${summary.gpu || "?"} +

Trace Analysis

+ ${summary.model || "?"} | ${summary.gpu || "?"}
- - - + + +
- ${activeTab === "summary" && html`<${SummaryPage} summary=${summary} />`} + ${activeTab === "dashboard" && html`<${Dashboard} summary=${summary} detail=${detail} batchList=${batchList} activeBatch=${activeBatch} setActiveBatch=${setActiveBatch} />`} ${activeTab === "batch" && html`
${batchList.map((b) => html` - + `)}
- ${detail ? html` - <${KernelTable} kt=${detail.kernel_table} batch=${activeBatch} stage=${activeStage} /> - <${CategoryChart} kt=${detail.kernel_table} /> - <${OverlapPanel} ov=${detail.overlap} fu=${detail.fuse} /> - ` : html`
Loading batch detail…
`} + ${detail ? html`<${KernelTable} kt=${detail.kernel_table} batch=${activeBatch} stage=${activeStage} />` : html`
Loading…
`} `} - ${activeTab === "hints" && html`<${HintsPage} hints=${hints} />`} + ${activeTab === "hints" && html`<${HintsPage} hints=${hints} detail=${detail} />`}
`; } -/* ── Summary Overview ── */ +/* ═══════════════════════════════════════════════════════════════════════ + DASHBOARD + ═══════════════════════════════════════════════════════════════════════ */ + +function Dashboard({ summary, detail, batchList, activeBatch, setActiveBatch }) { + if (!detail) return html`
Loading dashboard…
`; + const kt = detail.kernel_table; + if (!kt) return null; + const kernels = kt.kernels || []; + + // Compute stats + const top = kernels[0] || {}; + const cats = {}; + let mfuVals = [], totalDur = kt.total_gpu_time_s || 0; + for (const k of kernels) { + cats[k.category] = (cats[k.category] || 0) + (k.time_pct || 0); + if (k.mfu != null && k.mfu > 0) mfuVals.push(k.mfu); + if (k.tflops_actual != null && k.tflops_actual > 0) mfuVals.push(k.tflops_actual / (k.tflops_theoretical || 192) * 100); + } + const avgMfu = mfuVals.length ? (mfuVals.reduce((a, b) => a + b, 0) / mfuVals.length).toFixed(1) : null; + const computePct = kernels.filter(k => k.bound === "compute").reduce((s, k) => s + (k.time_pct || 0), 0); + const memoryPct = kernels.filter(k => k.bound === "memory").reduce((s, k) => s + (k.time_pct || 0), 0); + const unknownBound = 100 - computePct - memoryPct; + + const ov = detail.overlap || {}; + const cudaGraphOk = (ov.summary || {}).cuda_graph_effective; + const gapCount = (ov.gaps || []).length; + + // Category colors + const catColors = { Reduce: "#c0392b", GEMM: "#d35400", ElementWise: "#e67e22", MoE: "#27ae60", + NCCL: "#e74c3c", Attention: "#8e44ad", Norm: "#2980b9", Indexing: "#16a085", Memory: "#f1c40f", + Quantization: "#2c3e50", Other: "#7f8c8d", Transform: "#2ecc71", Activation: "#e91e63" }; + + const sortedCats = Object.entries(cats).sort((a, b) => b[1] - a[1]); + const totalPct = sortedCats.reduce((s, [, v]) => s + v, 0); + // Build conic-gradient stops for the donut + const donutStops = []; + let acc = 0; + for (const [cat, pct] of sortedCats) { + donutStops.push(`${catColors[cat] || "#95a5a6"} ${acc}% ${acc + pct}%`); + acc += pct; + } -function SummaryPage({ summary }) { - const batches = summary.batches || []; return html` -
-

Batch Summary

- - - - - - ${batches.map((b) => html` - - - - - - - - - `)} - -
BatchStageTop Kernel%KernelsMFU Avg
${b.batch_size}${b.stage}${(b.top_kernel || "").slice(0, 70)}${(b.top_kernel_pct || 0).toFixed(1)}%${b.kernel_count}${b.mfu_avg != null ? b.mfu_avg.toFixed(1) + "%" : "-"}
+
+ ${/* Row 1: Quick stats */""} +
+
+
${totalDur.toFixed(2)}s
+
Total GPU Time
+
+
+
${(top.time_pct || 0).toFixed(1)}%
+
Top Bottleneck
+
${(top.kernel_name || "").slice(0, 40)}
+
+
+
${avgMfu != null ? avgMfu + "%" : "—"}
+
Avg MFU (BF16)
+
+
+
${cudaGraphOk ? "ON" : "OFF"}
+
CUDA Graph
+
+
+ + ${/* Row 2: Category donut + Bound breakdown + Bottleneck detail */""} +
+
+

GPU Time by Category

+
+
+
+ ${kernels.length} + kernels +
+
+
+ ${sortedCats.slice(0, 8).map(([cat, pct]) => html` +
+ + ${cat} + ${pct.toFixed(1)}% +
+ `)} +
+
+
+ +
+

Bottleneck Detail

+
+
#1
+
+
${top.kernel_name || "?"}
+
+ Category: ${top.category || "?"} | + Op: ${top.op_type || "?"} | + Count: ${top.count || 0} +
+
+ Layer: ${top.model_layer || "unknown"} | + Bound: ${top.bound || "unknown"} | + MFU: ${top.mfu != null ? top.mfu.toFixed(1) + "%" : "—"} +
+
+
+
+
+
+ +

Compute vs Memory Bound

+
+
+ Compute-bound +
+ ${computePct.toFixed(1)}% +
+
+ Memory-bound +
+ ${memoryPct.toFixed(1)}% +
+
+ Unknown +
+ ${unknownBound.toFixed(1)}% +
+
+ +

Overlap

+

${gapCount} idle gaps detected. ${cudaGraphOk ? "CUDA Graph is active — gaps are minimal." : "CUDA Graph is OFF — explore enabling it."}

+
+
+ + ${/* Row 3: Top kernels quick preview */""} +
+

Top Kernels

+ + + + ${kernels.slice(0, 10).map((k) => html` + + + + + + + + + + + `)} + +
#%CategoryKernelCountAvg μsMFUBound
${k.rank}
${(k.time_pct || 0).toFixed(1)}%
${k.category || "?"}${(k.kernel_name || "").slice(0, 55)}${k.count}${(k.avg_dur_us || 0).toFixed(1)}${k.mfu != null ? k.mfu.toFixed(1) + "%" : "—"}${k.bound || "?"}
+
`; } -/* ── Kernel Hotspot Table ── */ +/* ═══════════════════════════════════════════════════════════════════════ + KERNEL TABLE (full, searchable) + ═══════════════════════════════════════════════════════════════════════ */ function KernelTable({ kt, batch, stage }) { if (!kt) return null; const kernels = kt.kernels || []; const totalTime = kt.total_gpu_time_s || 1; + const [search, setSearch] = useState(""); + const [catFilter, setCatFilter] = useState("all"); + + const categories = [...new Set(kernels.map((k) => k.category || "Other"))]; + const filtered = kernels.filter((k) => { + if (catFilter !== "all" && k.category !== catFilter) return false; + if (search && !k.kernel_name.toLowerCase().includes(search.toLowerCase())) return false; + return true; + }); return html`

Kernel Hotspots — BS=${batch} ${stage} (${kernels.length} unique, ${totalTime.toFixed(1)}s GPU)

+
+ setSearch(e.target.value)} /> + + ${filtered.length} of ${kernels.length} kernels +
- - - - - - ${kernels.slice(0, 25).map((k) => html` - - - - - - - - - - - - - `)} - -
#%CategoryOpLayerMFUBoundCountAvg μsKernel
${k.rank}
${(k.time_pct || 0).toFixed(1)}%
${k.category || "?"}${k.op_type || "?"}${k.model_layer || "-"}${k.mfu != null ? k.mfu.toFixed(1) + "%" : "-"}${k.bound || "-"}${k.count}${(k.avg_dur_us || 0).toFixed(1)}${(k.kernel_name || "").slice(0, 60)}
+ + + + ${filtered.slice(0, 100).map((k) => html` + + + + + + + + + + + + + + `)} + +
#%CategoryOpLayerCountAvg μsMFUBoundConfKernel
${k.rank}
${(k.time_pct || 0).toFixed(1)}%
${k.category || "?"}${k.op_type || "?"}${k.model_layer || "—"}${k.count}${(k.avg_dur_us || 0).toFixed(1)}${k.mfu != null ? k.mfu.toFixed(1) + "%" : "—"}${k.bound || "—"}${k.confidence || "?"}${(k.kernel_name || "").slice(0, 60)}
`; } -/* ── Category Breakdown ── */ +/* ═══════════════════════════════════════════════════════════════════════ + HINTS PAGE + ═══════════════════════════════════════════════════════════════════════ */ -function CategoryChart({ kt }) { - if (!kt) return null; - const kernels = kt.kernels || []; - const cats = {}; - for (const k of kernels) { - cats[k.category] = (cats[k.category] || 0) + (k.time_pct || 0); - } - const sorted = Object.entries(cats).sort((a, b) => b[1] - a[1]); - const colors = ["#c0392b","#d35400","#e67e22","#27ae60","#2980b9","#8e44ad","#16a085","#7f8c8d","#2c3e50","#e91e63"]; +function HintsPage({ hints, detail }) { + const kt = detail ? detail.kernel_table : null; + const fuse = detail ? detail.fuse : null; + const fuseMatches = (fuse && fuse.matches) || []; return html` -
-

Category Breakdown

-
- ${sorted.map(([cat, pct], i) => html` -
- ${cat} -
-
-
- ${pct.toFixed(1)}% + ${kt && html` + <${BottleneckAnalysis} kt=${kt} /> + `} + + ${fuseMatches.length > 0 && html` +
+

Fuse Opportunities (${fuseMatches.length})

+ ${fuseMatches.map((m) => html` +
+ ${m.pattern} + ${m.confidence} + ~${m.estimated_saving_us}μs saving +

${m.suggestion}

`)}
-
- `; -} - -/* ── Overlap + Fuse ── */ - -function OverlapPanel({ ov, fu }) { - if (!ov) return null; - const sum = ov.summary || {}; - const gaps = ov.gaps || []; - const high = gaps.filter((g) => g.severity === "high").length; - const medium = gaps.filter((g) => g.severity === "medium").length; - const fuseMatches = (fu && fu.matches) || []; + `} - return html` -
-

Overlap & Fuse

-

- ${gaps.length} GPU idle gaps (${high} high, ${medium} medium) - — total: ${(sum.total_gap_us / 1000).toFixed(1)}ms - (${sum.total_gap_pct || 0}% of GPU time) -

-

CUDA Graph effective: ${sum.cuda_graph_effective ? "YES" : "NO"}

- ${fuseMatches.length > 0 && html` -

Fuse Pattern Matches

- ${fuseMatches.map((m) => html` -
- ${m.pattern} (${m.confidence}) -

${m.suggestion}

+ ${hints && hints.status !== "skipped" && html` +
+

AI Optimization Hints

+ ${(hints.suggestions || []).map((s) => html` +
+ ${s.title} + ${s.difficulty} + Est. saving: ${s.estimated_saving_pct}% +

${s.what_to_change}

+

${s.why} | Type: ${s.category}

`)} - `} -
+
+ `} + + ${(hints && hints.status === "skipped" && fuseMatches.length === 0 && !kt) && html` +

Optimization Hints

No hints or fuse matches available yet.

+ `} `; } -/* ── Hints Page ── */ - -function HintsPage({ hints }) { - if (!hints) return html`

No hints generated yet.

`; +function BottleneckAnalysis({ kt }) { + if (!kt) return null; + const kernels = kt.kernels || []; + const top = kernels[0]; + const top3 = kernels.slice(0, 3); - const b = hints.bottleneck || {}; - const suggestions = hints.suggestions || []; - const surprises = hints.surprises || []; - - if (hints.status === "skipped") { - return html`

Optimization Hints

Hints generation skipped (${hints.reason || "not wired"}).

`; - } + const computeBoundPct = kernels.filter(k => k.bound === "compute").reduce((s, k) => s + (k.time_pct || 0), 0); + const suggestions = []; + if (computeBoundPct < 30) suggestions.push("Most kernels are memory-bound — focus on kernel fusion to reduce memory traffic."); + if ((top.time_pct || 0) > 50) suggestions.push(`"${(top.kernel_name || "").slice(0, 40)}" dominates at ${(top.time_pct || 0).toFixed(1)}%. Consider optimizing or replacing this kernel.`); + if (suggestions.length === 0) suggestions.push("GPU time is spread across many kernels. Look for fusion opportunities in the table below."); return html`
-

Optimization Hints

-
-

Biggest Bottleneck

-

${b.kernel_or_pattern || "?"} — ${b.reason || ""} (impact: ${b.impact_pct || 0}%)

+

Bottleneck Analysis

+
+ ${top3.map((k, i) => html` +
+ #${i + 1} + ${(k.time_pct || 0).toFixed(1)}% + ${(k.kernel_name || "").slice(0, 60)} + ${k.category || "?"} +
+ `)}
- ${suggestions.length > 0 && html` -
-

Suggestions

- ${suggestions.map((s) => html` -
- ${s.title} ${s.difficulty} -

${s.what_to_change}

-

Why: ${s.why} | Est. saving: ${s.estimated_saving_pct}% | Type: ${s.category}

-
- `)} -
- `} - ${surprises.length > 0 && html` -
-

Surprises

- ${surprises.map((s) => html`

${s}

`)} -
- `} + ${suggestions.map((s) => html`

${s}

`)}
`; } diff --git a/metainfer/tasks/sglang_trace_analyze/static/sa.css b/metainfer/tasks/sglang_trace_analyze/static/sa.css index ca23dcb2..a665e874 100644 --- a/metainfer/tasks/sglang_trace_analyze/static/sa.css +++ b/metainfer/tasks/sglang_trace_analyze/static/sa.css @@ -1,25 +1,94 @@ +/* sglang_trace_analyze — dashboard + kernel table styles */ .sa-detail { padding: 12px 16px; color: #e0e0e0; font-family: system-ui, sans-serif; } -.sa-header { margin-bottom: 12px; } +.sa-header { margin-bottom: 8px; } .sa-header h2 { margin: 0 0 2px; color: #fff; font-size: 18px; } .sa-meta { color: #888; font-size: 12px; } .sa-loading,.sa-error,.sa-empty { padding: 32px; text-align: center; color: #888; } .sa-error { color: #e74c3c; } /* Tabs */ -.sa-tabs { display: flex; gap: 2px; margin-bottom: 12px; border-bottom: 2px solid #333; } +.sa-tabs { display: flex; gap: 2px; margin-bottom: 10px; border-bottom: 2px solid #333; } .sa-tab-btn { padding: 6px 16px; border: none; border-radius: 4px 4px 0 0; background: transparent; color: #999; cursor: pointer; font-size: 13px; } .sa-tab-btn.active-tab { background: #2a2a2a; color: #4a90d9; font-weight: 600; } -.sa-batch-tabs { display: flex; gap: 6px; margin-bottom: 12px; } +.sa-batch-tabs { display: flex; gap: 6px; margin-bottom: 10px; } .sa-tab { padding: 4px 12px; border: 1px solid #444; border-radius: 4px; background: #2a2a2a; color: #ccc; cursor: pointer; font-size: 12px; } .sa-tab.active { background: #4a90d9; color: #fff; border-color: #4a90d9; } /* Panels */ .sa-panel { background: #1e1e1e; border: 1px solid #333; border-radius: 6px; - padding: 14px; margin-bottom: 12px; } -.sa-panel h3 { margin: 0 0 10px; color: #ddd; font-size: 14px; } + padding: 14px; margin-bottom: 10px; } +.sa-panel h3 { margin: 0 0 8px; color: #ddd; font-size: 13px; } .sa-note { color: #888; font-size: 11px; margin: 4px 0; } +.ml8 { margin-left: 8px; } + +/* Stat cards */ +.sa-stats-row { display: grid; grid-template-columns: repeat(4, 1fr); gap: 10px; margin-bottom: 10px; } +.sa-stat-card { background: #1e1e1e; border: 1px solid #333; border-radius: 6px; + padding: 12px; text-align: center; } +.sa-stat-value { font-size: 24px; font-weight: 700; color: #fff; line-height: 1.2; } +.sa-stat-label { font-size: 11px; color: #888; margin-top: 4px; } +.sa-stat-sub { font-size: 10px; color: #666; margin-top: 2px; + overflow: hidden; text-overflow: ellipsis; white-space: nowrap; max-width: 180px; } +.sa-stat-warn { border-color: #c0392b; } +.sa-stat-warn .sa-stat-value { color: #e74c3c; } +.sa-stat-ok { border-color: #27ae60; } +.sa-stat-ok .sa-stat-value { color: #2ecc71; } + +/* Grid */ +.sa-grid-2col { display: grid; grid-template-columns: 380px 1fr; gap: 10px; margin-bottom: 10px; } + +/* Donut chart */ +.sa-donut-wrap { display: flex; align-items: center; gap: 16px; } +.sa-donut { width: 140px; height: 140px; border-radius: 50%; position: relative; flex-shrink: 0; } +.sa-donut-hole { position: absolute; top: 28px; left: 28px; right: 28px; bottom: 28px; + background: #1e1e1e; border-radius: 50%; display: flex; flex-direction: column; + align-items: center; justify-content: center; } +.sa-donut-val { font-size: 22px; font-weight: 700; color: #fff; } +.sa-donut-lbl { font-size: 10px; color: #888; } +.sa-donut-legend { flex: 1; min-width: 0; } +.sa-legend-item { display: flex; align-items: center; gap: 6px; margin-bottom: 3px; } +.sa-legend-swatch { width: 10px; height: 10px; border-radius: 2px; flex-shrink: 0; } +.sa-legend-name { font-size: 11px; color: #ccc; overflow: hidden; text-overflow: ellipsis; white-space: nowrap; } +.sa-legend-pct { font-size: 11px; color: #aaa; margin-left: auto; font-variant-numeric: tabular-nums; } + +/* Bottleneck */ +.sa-bottleneck { display: flex; gap: 12px; } +.sa-bn-rank { font-size: 32px; font-weight: 800; color: #c0392b; line-height: 1; flex-shrink: 0; } +.sa-bn-info { flex: 1; min-width: 0; } +.sa-bn-name { font-size: 12px; font-family: monospace; color: #e74c3c; margin-bottom: 4px; + word-break: break-all; } +.sa-bn-meta { font-size: 11px; color: #888; margin-bottom: 2px; } +.sa-bn-meta strong { color: #ccc; } +.sa-bn-bar-wrap { background: #2a2a2a; border-radius: 3px; height: 20px; overflow: hidden; margin-top: 6px; } +.sa-bn-bar { height: 100%; background: #c0392b; border-radius: 3px; min-width: 2px; } + +/* Bound bars */ +.sa-bound-bars { } +.sa-bound-row { display: flex; align-items: center; gap: 8px; margin-bottom: 5px; } +.sa-bound-label { width: 100px; font-size: 11px; color: #ccc; text-align: right; flex-shrink: 0; } +.sa-bound-bar-bg { flex: 1; background: #2a2a2a; border-radius: 3px; height: 14px; overflow: hidden; } +.sa-bound-bar { height: 100%; border-radius: 3px; min-width: 2px; } +.sa-bb-compute { background: #2ecc71; } +.sa-bb-memory { background: #e67e22; } +.sa-bb-unknown { background: #555; } +.sa-bound-pct { width: 45px; font-size: 11px; color: #aaa; font-variant-numeric: tabular-nums; } + +/* Bottleneck list (Hints page) */ +.sa-bn-list { margin-bottom: 10px; } +.sa-bn-row { display: flex; align-items: center; gap: 8px; padding: 4px 0; border-bottom: 1px solid #2a2a2a; } +.sa-bn-rank-sm { font-size: 12px; font-weight: 700; color: #888; min-width: 24px; } +.sa-bn-pct { font-size: 13px; font-weight: 600; color: #e74c3c; min-width: 48px; font-variant-numeric: tabular-nums; } +.sa-bn-name-sm { font-size: 11px; font-family: monospace; color: #ccc; flex: 1; overflow: hidden; text-overflow: ellipsis; white-space: nowrap; } + +/* Filter bar */ +.sa-filters { display: flex; gap: 8px; align-items: center; margin-bottom: 8px; } +.sa-search { background: #2a2a2a; border: 1px solid #444; border-radius: 4px; padding: 4px 8px; + color: #ddd; font-size: 12px; width: 220px; } +.sa-select { background: #2a2a2a; border: 1px solid #444; border-radius: 4px; padding: 4px 8px; + color: #ddd; font-size: 12px; } +.sa-filter-count { font-size: 11px; color: #888; } /* Table */ .sa-table-wrap { overflow-x: auto; } @@ -29,7 +98,7 @@ .sa-table td { padding: 3px 6px; border-bottom: 1px solid #2a2a2a; vertical-align: middle; } .sa-num { text-align: right; font-variant-numeric: tabular-nums; color: #aaa; } .sa-sm { font-size: 10px; color: #888; } -.sa-kernel-name { max-width: 280px; overflow: hidden; text-overflow: ellipsis; +.sa-kernel-name { max-width: 300px; overflow: hidden; text-overflow: ellipsis; white-space: nowrap; font-family: monospace; font-size: 10px; color: #bbb; } .sa-pct { width: 100px; } .sa-bar-bg { position: relative; background: #2a2a2a; border-radius: 2px; @@ -41,25 +110,20 @@ .sa-cat { display: inline-block; padding: 1px 5px; border-radius: 2px; font-size: 10px; background: #333; color: #ccc; } -/* Category chart */ -.sa-cat-chart { display: flex; flex-direction: column; gap: 5px; } -.sa-cat-row { display: flex; align-items: center; gap: 8px; } -.sa-cat-label { width: 120px; font-size: 11px; color: #ccc; text-align: right; } -.sa-cat-bar-bg { flex: 1; background: #2a2a2a; border-radius: 2px; height: 16px; overflow: hidden; } -.sa-cat-bar { height: 100%; border-radius: 2px; min-width: 2px; } -.sa-cat-pct { width: 50px; font-size: 11px; color: #aaa; font-variant-numeric: tabular-nums; } +/* Confidence badges */ +.sa-conf { display: inline-block; padding: 1px 4px; border-radius: 2px; font-size: 10px; } +.sa-conf-high { background: #27ae60; color: #fff; } +.sa-conf-medium { background: #e67e22; color: #fff; } +.sa-conf-low { background: #c0392b; color: #fff; } /* Hints */ -.sa-hint-section { margin-top: 10px; } -.sa-hint-section h4 { margin: 0 0 6px; color: #ccc; font-size: 12px; } .sa-hint-card { background: #252525; border-left: 3px solid #4a90d9; padding: 8px 10px; margin-bottom: 8px; border-radius: 0 4px 4px 0; } +.sa-hint-card strong { color: #f1c40f; } .sa-difficulty { display: inline-block; padding: 1px 6px; border-radius: 3px; font-size: 10px; color: #fff; margin-left: 6px; } .sa-diff-low { background: #27ae60; } .sa-diff-medium { background: #e67e22; } .sa-diff-high { background: #c0392b; } - -/* Fuse cards */ -.sa-fuse-card { background: #252525; padding: 6px 8px; margin: 6px 0; border-radius: 4px; } -.sa-fuse-card strong { color: #f1c40f; } +.sa-suggestion { font-size: 12px; color: #f1c40f; background: #2a2a20; + border-left: 3px solid #f1c40f; padding: 6px 10px; margin: 6px 0; border-radius: 0 4px 4px 0; } From 8b29298ed720d12881397eab3a196a71a01b11a0 Mon Sep 17 00:00:00 2001 From: flyingdown Date: Thu, 6 Aug 2026 00:42:17 +0800 Subject: [PATCH 09/15] fix(sglang-trace-analyze): improve mapping confidence with kernel name tiering Kernel names that unambiguously identify op type (CK GEMM, flash_attn, fused_moe, NCCL, w8a8, cross_device_reduce, topk radix/gather) now get "high" confidence without requiring call stacks. Result: 82.4% of GPU time covered by high-confidence mappings. Low confidence restricted to generic elementwise/memory kernels (17.6%). Co-Authored-By: deepseek-v4-pro[1m] --- .../orchestrator/structure_mapper.py | 74 +++++++++++++++++-- 1 file changed, 69 insertions(+), 5 deletions(-) diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/structure_mapper.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/structure_mapper.py index 6eb4af0d..b900f9c1 100644 --- a/metainfer/tasks/sglang_trace_analyze/orchestrator/structure_mapper.py +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/structure_mapper.py @@ -53,12 +53,21 @@ def _map_one( layer = _infer_layer(call_stack, kernel_name, config, cpu_ops) op_type = _infer_op_type(kernel_name, call_stack, cpu_ops) - confidence = "high" + has_cpu_hint = bool(cpu_ops) + if not call_stack: - # Without call stacks, we use kernel name + CPU op correlation - has_cpu_hint = bool(cpu_ops) - if has_cpu_hint and _is_ck_gemm(kernel_name): - confidence = "medium" # CK GEMM is unambiguous even without stack + # Confidence tiers without call stack: + # high: kernel name unambiguously identifies op type + # (CK GEMM, flash_attn, fused_moe, NCCL, w8a8, cross_device_reduce) + # medium: CPU ops provide corroborating hint + # low: no useful signal from either source + name_clear = _kernel_name_is_clear(kernel_name, op_type) + cpu_confirms = _cpu_ops_confirm(kernel_name, cpu_ops, op_type) + + if name_clear: + confidence = "high" + elif cpu_confirms: + confidence = "medium" elif has_cpu_hint: confidence = "medium" else: @@ -81,6 +90,61 @@ def _is_ck_gemm(name: str) -> bool: return name.lower().startswith("cijk_") +def _kernel_name_is_clear(kernel_name: str, op_type: str) -> bool: + """Does the kernel name unambiguously identify its op type?""" + name_lower = kernel_name.lower() + # CK GEMM: name encodes tile dims, very clear + if name_lower.startswith("cijk_"): + return True + # Flash attention / MLA kernels + if "flash_fwd" in name_lower or "flash_attn" in name_lower: + return True + # Fused MoE + if "fused_moe" in name_lower: + return True + # NCCL operations + if "nccl" in name_lower: + return True + # w8a8 GEMM kernels (INT8 quantized) + if "w8a8" in name_lower and "scaled_mm" in name_lower: + return True + # Custom allreduce (cross_device_reduce) + if "cross_device_reduce" in name_lower: + return True + # MHC pre/post kernels + if "mhc_pre" in name_lower or "mhc_post" in name_lower: + return True + # topk kernels + if "topk" in name_lower and ("radix" in name_lower or "gather" in name_lower or "find" in name_lower): + return True + return False + + +def _cpu_ops_confirm( + kernel_name: str, + cpu_ops: list | None, + op_type: str, +) -> bool: + """Do the correlated CPU ops confirm the kernel's op type?""" + if not cpu_ops: + return False + cpu_lower = " ".join(cpu_ops).lower() + + confirmations = { + "GEMM": ["aten::linear", "aten::addmm", "aten::matmul", "torch.compile"], + "Attention": ["flash_attn", "flash_fwd", "attention"], + "MoE": ["fused_moe", "moe", "experts"], + "Norm": ["rms_norm", "rmsnorm", "layer_norm", "layernorm"], + "NCCL": ["allreduce", "allgather", "all_reduce", "nccl"], + "Reduce": ["all_reduce", "reduce", "cross_device"], + "ElementWise": ["copy_", "add", "mul", "silu", "gelu", "reshape", "view"], + } + + patterns = confirmations.get(op_type, []) + return any(p in cpu_lower for p in patterns) + + + def _infer_layer( call_stack: str, kernel_name: str, From e5881ed47a41f679b57f4c0503fec553682f116b Mon Sep 17 00:00:00 2001 From: flyingdown Date: Thu, 6 Aug 2026 00:55:51 +0800 Subject: [PATCH 10/15] feat(sglang-trace-analyze): add TFLOPS, bandwidth, structure mapping, fuse panels MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Dashboard now includes six analysis sections: - TFLOPS & Bandwidth table: actual vs theoretical peak per kernel - Model Structure → Operator Mapping: layer↔kernel groupings with confidence distribution, showing which model layers produce which GPU operators - Fuse Opportunities: rule-based pattern matches with estimated savings - Inefficiency Radar: kernels with high time + low MFU ranked by waste - Roofline Analysis: ops/byte vs ridge point visualization - Category donut chart + Compute/Memory bound + Bottleneck detail Mapping data fetched via /mapping API, confidence stats shown inline. Co-Authored-By: deepseek-v4-pro[1m] --- .../sglang_trace_analyze/static/sa-detail.js | 230 +++++++++++++++++- .../tasks/sglang_trace_analyze/static/sa.css | 23 ++ 2 files changed, 249 insertions(+), 4 deletions(-) diff --git a/metainfer/tasks/sglang_trace_analyze/static/sa-detail.js b/metainfer/tasks/sglang_trace_analyze/static/sa-detail.js index 58a7f3f3..bfd8cbb8 100644 --- a/metainfer/tasks/sglang_trace_analyze/static/sa-detail.js +++ b/metainfer/tasks/sglang_trace_analyze/static/sa-detail.js @@ -12,6 +12,7 @@ export default function SADetail({ taskId }) { const [summary, setSummary] = useState(null); const [hints, setHints] = useState(null); const [detail, setDetail] = useState(null); + const [mapping, setMapping] = useState(null); const [activeTab, setActiveTab] = useState("dashboard"); const [activeBatch, setActiveBatch] = useState(null); const [activeStage, setActiveStage] = useState("decode"); @@ -22,8 +23,9 @@ export default function SADetail({ taskId }) { Promise.all([ fetch(`${API(taskId)}/summary`).then((r) => r.json()), fetch(`${API(taskId)}/hints`).then((r) => r.json()), + fetch(`${API(taskId)}/mapping`).then((r) => r.json()), ]) - .then(([s, h]) => { setSummary(s); setHints(h); setLoading(false); }) + .then(([s, h, m]) => { setSummary(s); setHints(h); setMapping(m); setLoading(false); }) .catch((e) => { setError(e.message); setLoading(false); }); }, [taskId]); @@ -57,7 +59,7 @@ export default function SADetail({ taskId }) {
- ${activeTab === "dashboard" && html`<${Dashboard} summary=${summary} detail=${detail} batchList=${batchList} activeBatch=${activeBatch} setActiveBatch=${setActiveBatch} />`} + ${activeTab === "dashboard" && html`<${Dashboard} summary=${summary} detail=${detail} mapping=${mapping} batchList=${batchList} activeBatch=${activeBatch} setActiveBatch=${setActiveBatch} />`} ${activeTab === "batch" && html`
${batchList.map((b) => html` @@ -75,7 +77,7 @@ export default function SADetail({ taskId }) { DASHBOARD ═══════════════════════════════════════════════════════════════════════ */ -function Dashboard({ summary, detail, batchList, activeBatch, setActiveBatch }) { +function Dashboard({ summary, detail, mapping, batchList, activeBatch, setActiveBatch }) { if (!detail) return html`
Loading dashboard…
`; const kt = detail.kernel_table; if (!kt) return null; @@ -206,7 +208,22 @@ function Dashboard({ summary, detail, batchList, activeBatch, setActiveBatch })
- ${/* Row 3: Top kernels quick preview */""} + ${/* Row 3: TFLOPS & Bandwidth + Structure Mapping */""} +
+ <${TflopsPanel} kernels=${kernels} gpu=${summary.gpu || "K100"} /> + <${StructureMappingPanel} mapping=${mapping} kernels=${kernels} /> +
+ + ${/* Row 4: Fuse + Mapping confidence */""} + <${FusePanel} detail=${detail} /> + + ${/* Row 5: Inefficiency radar + roofline */""} +
+ <${InefficiencyRadar} kernels=${kernels} /> + <${RooflinePanel} kernels=${kernels} gpu=${summary.gpu || "K100"} /> +
+ + ${/* Row 4: Top kernels quick preview */""}

Top Kernels

@@ -364,3 +381,208 @@ function BottleneckAnalysis({ kt }) { `; } + +/* ── Inefficiency Radar: high-time, low-MFU kernels ── */ + +function InefficiencyRadar({ kernels }) { + if (!kernels || !kernels.length) return null; + // Top kernels by (time_pct * (100 - mfu)) / 100 — high time, low efficiency + const inefficiency = kernels + .filter((k) => (k.time_pct || 0) > 0.03) + .map((k) => ({ + ...k, + waste: ((k.time_pct || 0) * (k.mfu != null ? Math.max(0, 100 - k.mfu) : 100)) / 100, + })) + .sort((a, b) => b.waste - a.waste); + + return html` +
+

Inefficiency Radar

+

Kernels with high GPU time and low MFU — biggest optimization potential.

+
+ + + ${inefficiency.slice(0, 8).map((k) => html` + + + + + + + + `)} + +
KernelTime%MFUWaste ScoreCategory
${(k.kernel_name || "").slice(0, 50)}${(k.time_pct || 0).toFixed(1)}%${k.mfu != null ? k.mfu.toFixed(1) + "%" : "—"}${k.waste.toFixed(1).replace(/^-/, "")}${k.category || "?"}
+
+ `; +} + +/* ── Roofline Analysis ── */ + +function RooflinePanel({ kernels, gpu }) { + if (!kernels || !kernels.length) return null; + + // GPU peaks + const peaks = { K100: { bf16: 192, bw: 700 }, A100_80G: { bf16: 312, bw: 2039 }, + H100: { bf16: 989, bw: 3350 }, B200: { bf16: 2250, bw: 8000 } }; + const pk = peaks[gpu] || peaks.K100; + const peakFlops = pk.bf16 * 1e12; // TFLOPS → FLOPS + const peakBw = pk.bw * 1e9; // GB/s → B/s + const ridgePoint = peakFlops / peakBw; // ops/byte at the ridge + + // Classify each kernel with valid data + const pts = kernels + .filter((k) => k.tflops_actual != null && k.tflops_actual > 0 && k.bandwidth_gb_s != null && k.bandwidth_gb_s > 0) + .map((k) => ({ + name: k.kernel_name, category: k.category, time_pct: k.time_pct, + flops: k.tflops_actual * 1e12, bw: k.bandwidth_gb_s * 1e9, + opsPerByte: (k.tflops_actual * 1e12) / (k.bandwidth_gb_s * 1e9), + bound: k.bound, rank: k.rank, + })); + + const computeBound = pts.filter((p) => p.bound === "compute").length; + const memoryBound = pts.filter((p) => p.bound === "memory").length; + + return html` +
+

Roofline Analysis

+

GPU: ${gpu} | Peak BF16: ${pk.bf16} TFLOPS | BW: ${pk.bw} GB/s | Ridge: ${ridgePoint.toFixed(0)} ops/byte

+

+ ${computeBound} compute-bound | + ${memoryBound} memory-bound + ${pts.length < 5 ? html` (${kernels.length - pts.length} kernels lack dims for roofline)` : ""} +

+
+ ${pts.slice(0, 12).map((p) => { + const barW = Math.min(Math.log10(Math.max(p.opsPerByte, 1)) / Math.log10(ridgePoint * 10) * 100, 100); + const onRidge = p.opsPerByte > ridgePoint; + return html` +
+ ${(p.name || "").slice(0, 40)} + +
+ + ${p.opsPerByte.toFixed(0)} op/B + ${onRidge ? "compute" : "memory"} +
+ `; + })} +
+

Ridge point: ${ridgePoint.toFixed(0)} ops/byte. Left of ridge = memory-bound. Right = compute-bound.

+
+ `; +} + +/* ── TFLOPS & Bandwidth Panel ── */ + +function TflopsPanel({ kernels, gpu }) { + if (!kernels || !kernels.length) return null; + const peaks = { K100: { bf16: 192, bw: 700 }, A100_80G: { bf16: 312, bw: 2039 }, + H100: { bf16: 989, bw: 3350 }, B200: { bf16: 2250, bw: 8000 } }; + const pk = peaks[gpu] || peaks.K100; + + // Kernels with actual TFLOPS data + const withData = kernels.filter((k) => k.tflops_actual != null && k.tflops_actual > 0); + const withBw = kernels.filter((k) => k.bandwidth_gb_s != null && k.bandwidth_gb_s > 0); + + return html` +
+

TFLOPS & Bandwidth

+

GPU: ${gpu} | Theoretical peak BF16: ${pk.bf16} TFLOPS | BW: ${pk.bw} GB/s

+

${withData.length}/${kernels.length} kernels have TFLOPS data (CK GEMM tile dims extracted from kernel names).

+ + + + ${kernels.filter(k => k.tflops_actual != null || k.bandwidth_gb_s != null).slice(0, 10).map((k) => html` + + + + + + + + + `)} + +
KernelTFLOPSPeak%BW GB/sBW%Bound
${(k.kernel_name || "").slice(0, 45)}${k.tflops_actual != null ? k.tflops_actual.toFixed(3) : "—"}${k.mfu != null ? k.mfu.toFixed(1) + "%" : "—"}${k.bandwidth_gb_s != null ? k.bandwidth_gb_s.toFixed(1) : "—"}${k.bandwidth_gb_s != null ? (k.bandwidth_gb_s / pk.bw * 100).toFixed(1) + "%" : "—"}${k.bound || "—"}
+
+ `; +} + +/* ── Model Structure → Operator Mapping Panel ── */ + +function StructureMappingPanel({ mapping, kernels }) { + if (!mapping || !mapping.entries) return html`

Model Structure Mapping

No mapping data available.

`; + + const entries = mapping.entries || []; + // Group by model_layer + const layerGroups = {}; + for (const e of entries) { + const layer = e.model_layer || "unknown"; + if (!layerGroups[layer]) layerGroups[layer] = { kernels: [], categories: {} }; + layerGroups[layer].kernels.push(e); + layerGroups[layer].categories[e.category] = (layerGroups[layer].categories[e.category] || 0) + 1; + } + + const layers = Object.entries(layerGroups).sort((a, b) => b[1].kernels.length - a[1].kernels.length); + + // Confidence stats + const confStats = { high: 0, medium: 0, low: 0 }; + for (const e of entries) { confStats[e.confidence || "low"]++; } + const total = entries.length || 1; + + return html` +
+

Model Structure → Operator Mapping

+

${entries.length} kernel↔layer mappings | + high ${confStats.high} (${(confStats.high/total*100).toFixed(0)}%) + med ${confStats.medium} (${(confStats.medium/total*100).toFixed(0)}%) + low ${confStats.low} (${(confStats.low/total*100).toFixed(0)}%) +

+
+ ${layers.slice(0, 10).map(([layer, group]) => html` +
+ ${layer} + ${group.kernels.length} kernels + + ${Object.entries(group.categories).slice(0, 4).map(([cat, n]) => html` + ${cat}×${n} + `)} + +
+ `)} +
+
+ `; +} + +/* ── Fuse Opportunities Panel ── */ + +function FusePanel({ detail }) { + const fuse = detail ? detail.fuse : null; + const matches = fuse ? (fuse.matches || []) : []; + + if (matches.length === 0) return html` +
+

Fuse Opportunities

+

No fuse pattern matches found in rule engine. Try enabling LLM hints for AI-generated suggestions.

+
+ `; + + return html` +
+

Fuse Opportunities (${matches.length})

+ ${matches.map((m) => html` +
+
+ ${m.pattern} + ${m.confidence} + ~${m.estimated_saving_us}μs estimated saving +
+

${m.suggestion}

+

Kernels: ${(m.kernels || []).join(" → ")}

+
+ `)} +
+ `; +} diff --git a/metainfer/tasks/sglang_trace_analyze/static/sa.css b/metainfer/tasks/sglang_trace_analyze/static/sa.css index a665e874..f905f96c 100644 --- a/metainfer/tasks/sglang_trace_analyze/static/sa.css +++ b/metainfer/tasks/sglang_trace_analyze/static/sa.css @@ -127,3 +127,26 @@ .sa-diff-high { background: #c0392b; } .sa-suggestion { font-size: 12px; color: #f1c40f; background: #2a2a20; border-left: 3px solid #f1c40f; padding: 6px 10px; margin: 6px 0; border-radius: 0 4px 4px 0; } + +/* Roofline */ +.sa-roofline-bars { margin-top: 8px; } +.sa-rf-row { display: flex; align-items: center; gap: 6px; margin-bottom: 4px; } +.sa-rf-name { font-size: 10px; font-family: monospace; color: #bbb; width: 160px; overflow: hidden; text-overflow: ellipsis; white-space: nowrap; flex-shrink: 0; } +.sa-rf-bar-wrap { flex: 1; background: #2a2a2a; border-radius: 2px; height: 12px; overflow: hidden; } +.sa-rf-bar { height: 100%; border-radius: 2px; min-width: 2px; } +.sa-rf-compute { background: #2ecc71; } +.sa-rf-memory { background: #e67e22; } +.sa-rf-val { font-size: 10px; color: #888; width: 65px; text-align: right; font-variant-numeric: tabular-nums; flex-shrink: 0; } +.sa-rf-bound { font-size: 10px; color: #666; width: 60px; flex-shrink: 0; } + +/* Mapping grid */ +.sa-mapping-grid { margin-top: 8px; } +.sa-mapping-row { display: flex; align-items: center; gap: 8px; padding: 3px 0; border-bottom: 1px solid #2a2a2a; } +.sa-mapping-layer { font-size: 11px; color: #ccc; min-width: 140px; font-family: monospace; } +.sa-mapping-count { font-size: 10px; color: #888; min-width: 60px; } +.sa-mapping-cats { display: flex; gap: 4px; flex-wrap: wrap; } + +/* Fuse cards */ +.sa-fuse-card { background: #252525; border-left: 3px solid #e67e22; padding: 8px 10px; margin-bottom: 6px; border-radius: 0 4px 4px 0; } +.sa-fuse-header { display: flex; align-items: center; gap: 8px; margin-bottom: 4px; } +.sa-fuse-header strong { color: #f1c40f; } From 8f7aa3c069557fe854a0c2d1994c1bb1ea366c41 Mon Sep 17 00:00:00 2001 From: flyingdown Date: Thu, 6 Aug 2026 01:02:41 +0800 Subject: [PATCH 11/15] fix(sglang-trace-analyze): prefer formal traces, detect CUDA Graph from filename - _find_trace_dir: search formal traces (bs_N/timestamp/) before mapping fallback, so ANALYZE uses CUDA Graph ON traces when available - Detect CUDA Graph from trace filename (_graph_ vs _nograph_) instead of relying on gap count heuristic in overlap detector Result: Dashboard shows CUDA Graph: ON (green) when formal traces are used, OFF (red) only for mapping-only runs. Co-Authored-By: deepseek-v4-pro[1m] --- .../orchestrator/pipeline.py | 69 +++++++++++++++---- 1 file changed, 55 insertions(+), 14 deletions(-) diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/pipeline.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/pipeline.py index cdce8387..10cac8e2 100644 --- a/metainfer/tasks/sglang_trace_analyze/orchestrator/pipeline.py +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/pipeline.py @@ -331,20 +331,11 @@ def _run_analyze(self) -> bool: any_ok = False for bs in self.batch_sizes: for stage in self.stages: - # Try formal traces first, fall back to mapping traces - trace_dir = self.workspace_dir / "traces" / f"bs_{bs}" / stage - if not trace_dir.exists(): - # Fallback: look for mapping trace subdir - map_base = self.workspace_dir / "traces" / "mapping" - if map_base.exists(): - ts_dirs = sorted(map_base.glob("*/")) # timestamp subdirs - if ts_dirs: - trace_dir = ts_dirs[0] - else: - trace_dir = map_base - else: - print(f"[pipeline] skipping bs_{bs}/{stage} — no trace dir") - continue + # Priority: formal CUDA Graph ON traces > mapping traces + trace_dir = self._find_trace_dir(bs, stage) + if trace_dir is None: + print(f"[pipeline] skipping bs_{bs}/{stage} — no trace dir") + continue traces = sorted(trace_dir.glob("*DECODE*.trace.json.gz")) if not traces: @@ -536,6 +527,50 @@ def _phase_is_done(self, phase: str) -> bool: return (self._analysis_dir / "summary.json").exists() return False + # ================================================================== # + # Trace discovery + # ================================================================== # + + def _find_trace_dir(self, bs: int, stage: str) -> Optional[Path]: + """Find the best trace directory for a (batch_size, stage) pair. + + Priority: formal traces (CUDA Graph ON, under ``bs_/``) > + mapping traces (CUDA Graph OFF, under ``mapping/``). + + sglang ``--profile-by-stage`` saves traces inside a timestamp + subdirectory, so we look there first. + """ + def _find_in(base: Path) -> Optional[Path]: + if not base.exists(): + return None + # Direct: bs_8/decode/*.trace.json.gz + direct = base / stage + if direct.exists(): + traces = list(direct.glob("*DECODE*.trace.json.gz")) + if traces: + return direct + # Timestamp subdir: bs_8//*.trace.json.gz + ts_dirs = sorted([d for d in base.iterdir() if d.is_dir()]) + for ts in ts_dirs: + traces = list(ts.glob("*DECODE*.trace.json.gz")) + if traces: + return ts + return None + + # 1. Formal traces + formal_base = self.workspace_dir / "traces" / f"bs_{bs}" + found = _find_in(formal_base) + if found: + return found + + # 2. Mapping traces + map_base = self.workspace_dir / "traces" / "mapping" + found = _find_in(map_base) + if found: + return found + + return None + # ================================================================== # # Helpers # ================================================================== # @@ -696,6 +731,12 @@ def _analyze_one( } overlap = build_overlap_report(trace_data, bs, stage) + # Detect CUDA Graph from trace filename: _graph_ = formal, _nograph_ = mapping + trace_name = str(trace_path) + if "_nograph_" in trace_name: + overlap["summary"]["cuda_graph_effective"] = False + elif "_graph_" in trace_name: + overlap["summary"]["cuda_graph_effective"] = True fuse = build_fuse_report(result_kernels, bs, stage) return {"kernel_table": kernel_table, "overlap": overlap, "fuse": fuse} From 9411e256850867c6955366bdd5e37072bed102bd Mon Sep 17 00:00:00 2001 From: flyingdown Date: Thu, 6 Aug 2026 01:14:04 +0800 Subject: [PATCH 12/15] feat(sglang-trace-analyze): add MFU distribution histogram and frequency analysis Dashboard additions: - MFU Distribution: histogram across 7 buckets (0-5%, 5-10%, ..., 90-100%) with avg/median stats, showing how efficiently the GPU is used - Top by Invocation Count: kernels ranked by call frequency, helping identify "death by a thousand cuts" patterns where many small invocations could be batched Co-Authored-By: deepseek-v4-pro[1m] --- .../sglang_trace_analyze/static/sa-detail.js | 85 ++++++++++++++++++- .../tasks/sglang_trace_analyze/static/sa.css | 8 ++ 2 files changed, 92 insertions(+), 1 deletion(-) diff --git a/metainfer/tasks/sglang_trace_analyze/static/sa-detail.js b/metainfer/tasks/sglang_trace_analyze/static/sa-detail.js index bfd8cbb8..c974ca91 100644 --- a/metainfer/tasks/sglang_trace_analyze/static/sa-detail.js +++ b/metainfer/tasks/sglang_trace_analyze/static/sa-detail.js @@ -223,7 +223,13 @@ function Dashboard({ summary, detail, mapping, batchList, activeBatch, setActive <${RooflinePanel} kernels=${kernels} gpu=${summary.gpu || "K100"} />
- ${/* Row 4: Top kernels quick preview */""} + ${/* Row 4: MFU Distribution + Frequency Analysis */""} +
+ <${MfuDistro} kernels=${kernels} gpu=${summary.gpu || "K100"} /> + <${FrequencyPanel} kernels=${kernels} /> +
+ + ${/* Row 5: Top kernels quick preview */""}

Top Kernels

@@ -473,6 +479,83 @@ function RooflinePanel({ kernels, gpu }) { `; } +/* ── MFU Distribution Histogram ── */ + +function MfuDistro({ kernels, gpu }) { + if (!kernels || !kernels.length) return null; + const peaks = { K100: { bf16: 192 }, A100_80G: { bf16: 312 }, H100: { bf16: 989 }, B200: { bf16: 2250 } }; + const pk = (peaks[gpu] || peaks.K100).bf16; + + // Compute MFU for ALL kernels from tflops_actual / theoretical + const mfuVals = kernels.map((k) => { + if (k.mfu != null) return k.mfu; + if (k.tflops_actual != null && k.tflops_actual > 0) return k.tflops_actual / pk * 100; + return null; + }).filter((v) => v != null); + + if (mfuVals.length === 0) return html`

MFU Distribution

No MFU data available (no Input Dims in trace).

`; + + const buckets = [0, 5, 10, 25, 50, 75, 90, 100]; + const labels = ["0-5%", "5-10%", "10-25%", "25-50%", "50-75%", "75-90%", "90-100%"]; + const hist = new Array(buckets.length - 1).fill(0); + for (const v of mfuVals) { + for (let i = buckets.length - 1; i >= 0; i--) { + if (v >= buckets[i]) { hist[i]++; break; } + } + } + + const maxN = Math.max(...hist, 1); + const avg = mfuVals.reduce((a, b) => a + b, 0) / mfuVals.length; + const median = mfuVals.sort((a, b) => a - b)[Math.floor(mfuVals.length / 2)]; + + return html` +
+

MFU Distribution

+

${mfuVals.length} kernels with TFLOPS data | avg=${avg.toFixed(1)}% | median=${median.toFixed(1)}%

+
+ ${hist.map((n, i) => html` +
+ ${labels[i]} +
+
+
+ ${n} +
+ `)} +
+
+ `; +} + +/* ── Frequency Analysis ── */ + +function FrequencyPanel({ kernels }) { + if (!kernels || !kernels.length) return null; + // Top kernels by call count + const byCount = [...kernels].sort((a, b) => (b.count || 0) - (a.count || 0)); + + return html` +
+

Top by Invocation Count

+

High invocation count kernels may indicate repeated small operations that could be batched.

+
+ + + ${byCount.slice(0, 10).map((k) => html` + + + + + + + + `)} + +
KernelCallsTime%Avg μsCategory
${(k.kernel_name || "").slice(0, 45)}${k.count}${(k.time_pct || 0).toFixed(1)}%${(k.avg_dur_us || 0).toFixed(1)}${k.category || "?"}
+
+ `; +} + /* ── TFLOPS & Bandwidth Panel ── */ function TflopsPanel({ kernels, gpu }) { diff --git a/metainfer/tasks/sglang_trace_analyze/static/sa.css b/metainfer/tasks/sglang_trace_analyze/static/sa.css index f905f96c..e4548aab 100644 --- a/metainfer/tasks/sglang_trace_analyze/static/sa.css +++ b/metainfer/tasks/sglang_trace_analyze/static/sa.css @@ -150,3 +150,11 @@ .sa-fuse-card { background: #252525; border-left: 3px solid #e67e22; padding: 8px 10px; margin-bottom: 6px; border-radius: 0 4px 4px 0; } .sa-fuse-header { display: flex; align-items: center; gap: 8px; margin-bottom: 4px; } .sa-fuse-header strong { color: #f1c40f; } + +/* MFU Histogram */ +.sa-hist { margin-top: 6px; } +.sa-hist-row { display: flex; align-items: center; gap: 8px; margin-bottom: 4px; } +.sa-hist-label { font-size: 10px; color: #888; width: 55px; text-align: right; flex-shrink: 0; } +.sa-hist-bar-bg { flex: 1; background: #2a2a2a; border-radius: 2px; height: 14px; overflow: hidden; } +.sa-hist-bar { height: 100%; background: #4a90d9; border-radius: 2px; min-width: 2px; } +.sa-hist-count { font-size: 10px; color: #aaa; width: 30px; text-align: right; font-variant-numeric: tabular-nums; flex-shrink: 0; } From 7a5f62199e1bfdd131fbe8cc0a3743f6311f9343 Mon Sep 17 00:00:00 2001 From: flyingdown Date: Thu, 6 Aug 2026 02:13:47 +0800 Subject: [PATCH 13/15] fix(sglang-trace-analyze): enrich formal traces with TFLOPS from mapping trace - _merge_mapping_tflops: cross-reference formal trace kernel table with mapping trace (CUDA Graph OFF) to fill in tflops_actual, mfu, bound, bandwidth_gb_s per kernel by name matching - Fix flops_calculator to preserve small TFLOPS values from CK GEMM tiles - _is_formal_trace: detect CUDA Graph status from trace filename Co-Authored-By: deepseek-v4-pro[1m] --- .../orchestrator/flops_calculator.py | 8 +- .../orchestrator/pipeline.py | 96 +++++++++++++++++++ 2 files changed, 100 insertions(+), 4 deletions(-) diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/flops_calculator.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/flops_calculator.py index 185392de..288d4a23 100644 --- a/metainfer/tasks/sglang_trace_analyze/orchestrator/flops_calculator.py +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/flops_calculator.py @@ -82,12 +82,12 @@ def calculate_mfu( crossover = float("inf") bound = "compute" if ops_per_byte > crossover else "memory" - k["tflops_actual"] = round(tflops_actual, 3) if tflops_actual > 0 else None + k["tflops_actual"] = round(tflops_actual, 6) if flops > 0 else None k["tflops_theoretical"] = theoretical_tflops - k["bandwidth_gb_s"] = round(bandwidth_gb_s, 1) if bandwidth_gb_s > 0 else None + k["bandwidth_gb_s"] = round(bandwidth_gb_s, 1) if bytes_moved > 0 else None k["bandwidth_theoretical"] = theoretical_bw - k["mfu"] = round(mfu, 1) if tflops_actual > 0 else None - k["bound"] = bound if (tflops_actual and tflops_actual > 0) else "unknown" + k["mfu"] = round(mfu, 3) if flops > 0 else None + k["bound"] = bound if (flops > 0 and bytes_moved > 0) else "unknown" k["flops_per_invocation"] = int(flops) return kernels diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/pipeline.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/pipeline.py index 10cac8e2..5bd68cc2 100644 --- a/metainfer/tasks/sglang_trace_analyze/orchestrator/pipeline.py +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/pipeline.py @@ -592,6 +592,98 @@ def _build_bench_config( "profile_steps": 5, } + def _is_formal_trace(self, trace_path) -> bool: + """Return True if this is a formal (CUDA Graph ON) trace.""" + return "_graph_" in str(trace_path) and "_nograph_" not in str(trace_path) + + def _merge_mapping_tflops( + self, result_kernels: list, bs: int, stage: str + ) -> None: + """Enrich formal trace kernel entries with TFLOPS/MFU/bound from + the mapping trace, which has per-kernel Input Dims. + + Matches kernels by name and overwrites tflops_actual, mfu, + bound, bandwidth_gb_s, and input_dims from the mapping trace. + """ + mapping_trace_dir = self._find_mapping_trace_dir() + if mapping_trace_dir is None: + print("[pipeline] no mapping trace to enrich TFLOPS data") + return + + traces = sorted(mapping_trace_dir.glob("*DECODE*.trace.json.gz")) + if not traces: + return + + from .trace_parser import parse_trace, aggregate_kernels + from .flops_calculator import calculate_mfu + + print(f"[pipeline] enriching TFLOPS from mapping trace") + map_data = parse_trace(str(traces[0])) + map_kernels = aggregate_kernels(map_data) + + # Build CPU op correlation for mapping trace too + events = map_data.get("traceEvents", []) + cpu_ops_by_corr = defaultdict(lambda: []) + kernel_by_corr = defaultdict(lambda: []) + for e in events: + cat = e.get("cat", "") + corr = (e.get("args") or {}).get("External id" if cat == "cpu_op" else "correlation") + if cat == "cpu_op" and corr: + cpu_ops_by_corr[corr].append(e.get("name", "")) + elif cat == "kernel" and corr: + kernel_by_corr[corr].append(e.get("name", "")) + + # Classify + calculate MFU + from .structure_mapper import _map_one as map_one + map_entries = [] + for k in map_kernels: + name = k["kernel_name"] + cpu_ops = set() + for corr, gpu_names in kernel_by_corr.items(): + if name in gpu_names: + for cn in cpu_ops_by_corr.get(corr, []): + cpu_ops.add(cn) + mapped = map_one(name, k.get("call_stack", ""), {}, list(cpu_ops)) + k.update(mapped) + map_entries.append(k) + + map_entries = calculate_mfu(map_entries, self.gpu_spec, batch_size=bs, dtype="bf16") + + # Build lookup by kernel name + map_lookup = {k["kernel_name"]: k for k in map_entries} + + enriched = 0 + for k in result_kernels: + name = k["kernel_name"] + if name in map_lookup: + src = map_lookup[name] + if src.get("tflops_actual") is not None: + k["tflops_actual"] = src["tflops_actual"] + k["mfu"] = src["mfu"] + k["bound"] = src["bound"] + k["bandwidth_gb_s"] = src["bandwidth_gb_s"] + k["input_dims"] = src.get("input_dims", []) + k["flops_per_invocation"] = src.get("flops_per_invocation", 0) + enriched += 1 + print(f"[pipeline] enriched {enriched}/{len(result_kernels)} kernels with TFLOPS from mapping trace") + + def _find_mapping_trace_dir(self) -> Optional[Path]: + """Find the mapping trace directory (CUDA Graph OFF).""" + map_base = self.workspace_dir / "traces" / "mapping" + if not map_base.exists(): + return None + # Check for timestamp subdirs first + ts_dirs = sorted([d for d in map_base.iterdir() if d.is_dir()]) + for ts in ts_dirs: + traces = list(ts.glob("*DECODE*.trace.json.gz")) + if traces: + return ts + # Direct + traces = list(map_base.glob("*DECODE*.trace.json.gz")) + if traces: + return map_base + return None + def _build_mapping(self, trace_path: Path) -> List[Dict[str, Any]]: """Parse a trace file and build kernel→model-structure mapping using trace_parser + structure_mapper with CPU op correlation.""" @@ -730,6 +822,10 @@ def _analyze_one( "kernels": result_kernels, } + # Enrich formal traces with TFLOPS from mapping trace + if self._is_formal_trace(trace_path): + self._merge_mapping_tflops(result_kernels, bs, stage) + overlap = build_overlap_report(trace_data, bs, stage) # Detect CUDA Graph from trace filename: _graph_ = formal, _nograph_ = mapping trace_name = str(trace_path) From 60d56a65e06cd4316f5416deb1bc748b70f05d1d Mon Sep 17 00:00:00 2001 From: flyingdown Date: Thu, 6 Aug 2026 03:12:22 +0800 Subject: [PATCH 14/15] feat(sglang-trace-analyze): add autogenerated Key Findings panel Auto-generates 5 insight cards from analysis data: - Dominant kernel alert (single kernel >30% GPU time) - CUDA Graph status assessment - Category concentration warning (>50% in one category) - Top-3 kernels summary with category + time_pct - MFU data availability note with actionable next step Cards use icon + color coding (red/yellow/green) for quick scanning. Co-Authored-By: deepseek-v4-pro[1m] --- .../sglang_trace_analyze/static/sa-detail.js | 104 +++++++++++++++++- .../tasks/sglang_trace_analyze/static/sa.css | 7 ++ 2 files changed, 109 insertions(+), 2 deletions(-) diff --git a/metainfer/tasks/sglang_trace_analyze/static/sa-detail.js b/metainfer/tasks/sglang_trace_analyze/static/sa-detail.js index c974ca91..06580a59 100644 --- a/metainfer/tasks/sglang_trace_analyze/static/sa-detail.js +++ b/metainfer/tasks/sglang_trace_analyze/static/sa-detail.js @@ -223,13 +223,16 @@ function Dashboard({ summary, detail, mapping, batchList, activeBatch, setActive <${RooflinePanel} kernels=${kernels} gpu=${summary.gpu || "K100"} />
- ${/* Row 4: MFU Distribution + Frequency Analysis */""} + ${/* Row 4: Key Findings */""} + <${KeyFindings} kernels=${kernels} kt=${kt} cudaGraph=${cudaGraphOk} /> + + ${/* Row 5: MFU Distribution + Frequency Analysis */""}
<${MfuDistro} kernels=${kernels} gpu=${summary.gpu || "K100"} /> <${FrequencyPanel} kernels=${kernels} />
- ${/* Row 5: Top kernels quick preview */""} + ${/* Row 6: Top kernels quick preview */""}

Top Kernels

@@ -479,6 +482,103 @@ function RooflinePanel({ kernels, gpu }) { `; } +/* ── Key Findings auto-summary ── */ + +function KeyFindings({ kernels, kt, cudaGraph }) { + if (!kernels || !kernels.length) return null; + + const total = kt.total_gpu_time_s || 0; + const top = kernels[0]; + const top3 = kernels.slice(0, 3); + + // Build findings from data + const findings = []; + + // 1. Dominant kernel + if ((top.time_pct || 0) > 30) { + findings.push({ + icon: "🔴", title: "Single kernel dominates", + text: `"${(top.kernel_name || "").slice(0, 45)}" consumes ${(top.time_pct || 0).toFixed(1)}% of GPU time alone. This is your primary optimization target.`, + }); + } else if ((top.time_pct || 0) > 15) { + findings.push({ + icon: "🟡", title: "Moderate hotspot", + text: `Top kernel "${(top.kernel_name || "").slice(0, 45)}" at ${(top.time_pct || 0).toFixed(1)}%. Consider fusion or replacement.`, + }); + } else { + findings.push({ + icon: "🟢", title: "Well-distributed workload", + text: "GPU time is spread across many kernels. Focus on fusion and reducing kernel launch overhead.", + }); + } + + // 2. CUDA Graph + if (cudaGraph) { + findings.push({ + icon: "🟢", title: "CUDA Graph active", + text: `Total GPU time: ${total.toFixed(2)}s with CUDA Graph. Kernel launch overhead is minimized.`, + }); + } else { + findings.push({ + icon: "🔴", title: "CUDA Graph disabled", + text: "Enable CUDA Graph to reduce kernel launch overhead and CPU-GPU synchronization. Expected 3-5x speedup on decode.", + }); + } + + // 3. Category concentration + const cats = {}; + for (const k of kernels) cats[k.category] = (cats[k.category] || 0) + (k.time_pct || 0); + const topCat = Object.entries(cats).sort((a, b) => b[1] - a[1])[0]; + if (topCat && topCat[1] > 50) { + findings.push({ + icon: "🔴", title: `Category "${topCat[0]}" dominates at ${topCat[1].toFixed(0)}%`, + text: topCat[0] === "Reduce" ? "TP allreduce is the bottleneck. Consider communication-computation overlap or reducing TP degree." : + topCat[0] === "GEMM" ? "GEMM is the bottleneck. Explore quantization (FP8/INT8) or faster GEMM backends." : + `Focus optimization efforts on ${topCat[0]} operations.`, + }); + } + + // 4. Top 3 summary + const top3Summary = top3.map((k, i) => + `#${i + 1} ${(k.category || "?").slice(0, 10)} ${(k.time_pct || 0).toFixed(1)}%` + ).join(" | "); + findings.push({ + icon: "📊", title: "Top 3 kernels", + text: top3Summary, + }); + + // 5. MFU note + const withMfu = kernels.filter((k) => k.mfu != null && k.mfu > 0); + if (withMfu.length === 0) { + findings.push({ + icon: "💡", title: "No MFU data available", + text: "Profiler was run without record_shapes=True. Enable it to get per-kernel TFLOPS and MFU analysis.", + }); + } else if (withMfu.length < 10) { + findings.push({ + icon: "💡", title: `MFU data available for ${withMfu.length} kernels`, + text: "Limited TFLOPS data (only CK GEMM tiles). Enable record_shapes=True for full MFU coverage.", + }); + } + + return html` +
+

Key Findings

+
+ ${findings.map((f) => html` +
+ ${f.icon} +
+ ${f.title} +

${f.text}

+
+
+ `)} +
+
+ `; +} + /* ── MFU Distribution Histogram ── */ function MfuDistro({ kernels, gpu }) { diff --git a/metainfer/tasks/sglang_trace_analyze/static/sa.css b/metainfer/tasks/sglang_trace_analyze/static/sa.css index e4548aab..d6c47d63 100644 --- a/metainfer/tasks/sglang_trace_analyze/static/sa.css +++ b/metainfer/tasks/sglang_trace_analyze/static/sa.css @@ -158,3 +158,10 @@ .sa-hist-bar-bg { flex: 1; background: #2a2a2a; border-radius: 2px; height: 14px; overflow: hidden; } .sa-hist-bar { height: 100%; background: #4a90d9; border-radius: 2px; min-width: 2px; } .sa-hist-count { font-size: 10px; color: #aaa; width: 30px; text-align: right; font-variant-numeric: tabular-nums; flex-shrink: 0; } + +/* Key Findings */ +.sa-findings { display: flex; flex-wrap: wrap; gap: 8px; } +.sa-finding-card { display: flex; gap: 8px; background: #252525; border-radius: 4px; padding: 8px 10px; flex: 1; min-width: 280px; max-width: calc(50% - 4px); } +.sa-finding-icon { font-size: 16px; flex-shrink: 0; line-height: 1.2; } +.sa-finding-body { min-width: 0; } +.sa-finding-body strong { font-size: 12px; color: #ddd; } From ecf3d4c2e1f34d87c926c55c61370c09ef9c53ce Mon Sep 17 00:00:00 2001 From: flyingdown Date: Thu, 6 Aug 2026 04:12:14 +0800 Subject: [PATCH 15/15] feat(sglang-trace-analyze): rule-based hints + executive summary banner - Replace LLM stub with rule-based hint generation from kernel table, overlap, and fuse data. Generates 2-5 suggestions with difficulty rating, estimated saving %, and category. - Add executive summary banner at top of Dashboard: one-line summary of CUDA Graph status, bottleneck, and top optimization opportunities. - Fix hint collection to use full kernel list (not just top-3) for accurate category aggregation. Co-Authored-By: deepseek-v4-pro[1m] --- .../orchestrator/pipeline.py | 127 ++++++++++++++++-- .../sglang_trace_analyze/static/sa-detail.js | 44 ++++++ .../tasks/sglang_trace_analyze/static/sa.css | 5 + 3 files changed, 166 insertions(+), 10 deletions(-) diff --git a/metainfer/tasks/sglang_trace_analyze/orchestrator/pipeline.py b/metainfer/tasks/sglang_trace_analyze/orchestrator/pipeline.py index 5bd68cc2..ee9ca3ef 100644 --- a/metainfer/tasks/sglang_trace_analyze/orchestrator/pipeline.py +++ b/metainfer/tasks/sglang_trace_analyze/orchestrator/pipeline.py @@ -402,7 +402,7 @@ def _run_hints(self) -> bool: self.store.write_iteration(n, rec.to_dict()) self.store.update_run(current_iteration=n) - # Collect summaries from all analyzed batches + # Collect full kernel tables + summaries from all analyzed batches kernel_summaries = [] overlap_summaries = [] fuse_summaries = [] @@ -414,10 +414,9 @@ def _run_hints(self) -> bool: ov = _load_json(out_dir / "overlap.json") fu = _load_json(out_dir / "fuse.json") if kt: - top3 = (kt.get("kernels", []) or [])[:3] kernel_summaries.append({ "batch_size": bs, "stage": stage, - "top_kernels": top3, + "all_kernels": kt.get("kernels", []), }) if ov: overlap_summaries.append(ov) @@ -843,16 +842,124 @@ def _llm_generate_hints( overlap_summaries: list, fuse_summaries: list, ) -> Dict[str, Any]: - """Generate optimization hints via LLM. + """Generate optimization hints from analysis data. - Placeholder — real impl calls SubAgentManager. + Uses rule-based analysis of kernel tables, overlap, and fuse results + to produce actionable optimization suggestions. """ + suggestions = [] + surprises = [] + + # Collect all kernels across batches/stages + all_kernels = [] + for ks in kernel_summaries: + for k in (ks.get("all_kernels") or []): + all_kernels.append(k) + + if not all_kernels: + return { + "bottleneck": {"kernel_or_pattern": "unknown", "reason": "no data", "impact_pct": 0}, + "suggestions": [], "surprises": [], + "status": "generated", + } + + top = all_kernels[0] if all_kernels else {} + top_name = top.get("kernel_name", "unknown") + top_cat = top.get("category", "Other") + top_pct = top.get("time_pct", 0) + + # Categorize kernels + cats = {} + for k in all_kernels: + c = k.get("category", "Other") + cats[c] = cats.get(c, 0) + (k.get("time_pct", 0) or 0) + + # 1. Dominant kernel analysis + if top_pct > 50: + suggestions.append({ + "title": f"Replace or optimize {top_cat} kernel", + "what_to_change": f"The \"{top_name[:40]}\" kernel dominates at {top_pct:.0f}% GPU time. Profile with Nsight Compute to identify micro-architectural bottlenecks, or replace with a vendor-optimized implementation.", + "why": f"Single kernel consuming >50% of GPU time is the highest-ROI optimization target.", + "estimated_saving_pct": round(top_pct * 0.3), + "difficulty": "high", + "category": "kernel_replace", + }) + elif top_pct > 20: + suggestions.append({ + "title": f"Profile {top_cat} kernel with Nsight", + "what_to_change": f"\"{top_name[:40]}\" at {top_pct:.0f}%. Use Nsight Compute to check occupancy, memory coalescing, and register pressure.", + "why": "Top kernel is a clear bottleneck. Micro-architectural optimization may yield 10-30% improvement.", + "estimated_saving_pct": round(top_pct * 0.2), + "difficulty": "medium", + "category": "kernel_replace", + }) + + # 2. Category-specific suggestions + reduce_pct = cats.get("Reduce", 0) + if reduce_pct > 30: + suggestions.append({ + "title": "Reduce TP allreduce overhead", + "what_to_change": "Custom allreduce consumes {:.0f}% GPU time. Try: (1) overlap communication with computation using separate CUDA streams, (2) reduce TP degree if memory permits, or (3) enable CUDA Graph to amortize launch overhead.".format(reduce_pct), + "why": "TP communication is the dominant cost. Even 10% reduction saves significant time.", + "estimated_saving_pct": round(reduce_pct * 0.25), + "difficulty": "medium", + "category": "overlap", + }) + + gemm_pct = cats.get("GEMM", 0) + if gemm_pct > 20: + suggestions.append({ + "title": "Quantize GEMMs to FP8 or INT8", + "what_to_change": "GEMM kernels consume {:.0f}% GPU time. Explore FP8 (w8a8) quantization for attention projections and FFN layers to double throughput.".format(gemm_pct), + "why": "GEMM is compute-heavy and benefits most from reduced precision.", + "estimated_saving_pct": round(gemm_pct * 0.4), + "difficulty": "medium", + "category": "config_tune", + }) + + element_pct = cats.get("ElementWise", 0) + if element_pct > 15: + suggestions.append({ + "title": "Fuse element-wise operations", + "what_to_change": f"Element-wise kernels consume {element_pct:.0f}% GPU time. These are memory-bound — fuse consecutive element-wise ops (add, mul, silu, norm) into single kernels to reduce memory traffic.", + "why": "Memory-bound element-wise ops benefit most from fusion, eliminating intermediate reads/writes.", + "estimated_saving_pct": round(element_pct * 0.4), + "difficulty": "low", + "category": "fuse", + }) + + # 3. CUDA Graph check (from overlap data) + any_cuda_graph = any( + s.get("summary", {}).get("cuda_graph_effective", False) + for s in overlap_summaries + ) + if not any_cuda_graph: + suggestions.append({ + "title": "Enable CUDA Graph for decode", + "what_to_change": "CUDA Graph is not active. Enable --cuda-graph-bs to capture and replay the decode graph. On K100 with DeepSeek V4, this typically yields 3-5x throughput improvement.", + "why": "Decode is launch-bound. CUDA Graph eliminates per-step kernel launch overhead.", + "estimated_saving_pct": 70, + "difficulty": "low", + "category": "config_tune", + }) + + # 4. Surprises + if reduce_pct < 5 and "NCCL" not in cats: + surprises.append("TP allreduce overhead is unexpectedly low — verify communication is actually happening (check TP degree).") + if gemm_pct > 50: + surprises.append("GEMM dominates at >50% — unexpected for a decode workload. Check if attention is correctly fused.") + + bottleneck = { + "kernel_or_pattern": top_name[:80] if top_name else "unknown", + "reason": f"Largest single consumer of GPU time at {top_pct:.1f}% (category: {top_cat})", + "impact_pct": round(top_pct), + } + return { - "bottleneck": {"kernel_or_pattern": "TBD", "reason": "", "impact_pct": 0}, - "suggestions": [], - "surprises": [], - "status": "skipped", - "reason": "LLM hints not yet wired", + "bottleneck": bottleneck, + "suggestions": suggestions[:5], + "surprises": surprises, + "status": "generated", } diff --git a/metainfer/tasks/sglang_trace_analyze/static/sa-detail.js b/metainfer/tasks/sglang_trace_analyze/static/sa-detail.js index 06580a59..bf6f8d66 100644 --- a/metainfer/tasks/sglang_trace_analyze/static/sa-detail.js +++ b/metainfer/tasks/sglang_trace_analyze/static/sa-detail.js @@ -116,8 +116,14 @@ function Dashboard({ summary, detail, mapping, batchList, activeBatch, setActive acc += pct; } + const summaryText = buildSummary(kernels, cudaGraphOk, top); + return html`
+
+ ${summaryText} +
+ ${/* Row 1: Quick stats */""}
@@ -482,6 +488,44 @@ function RooflinePanel({ kernels, gpu }) { `; } +/* ── Executive Summary builder ── */ + +function buildSummary(kernels, cudaGraph, top) { + if (!kernels || !kernels.length) return "No analysis data available."; + + const parts = []; + parts.push(cudaGraph ? "CUDA Graph ON" : "CUDA Graph OFF"); + + if (top && top.category) { + parts.push(`${top.category} is your bottleneck (${(top.time_pct || 0).toFixed(0)}%)`); + } + + // Find category insights + const cats = {}; + for (const k of kernels) cats[k.category] = (cats[k.category] || 0) + (k.time_pct || 0); + + const reducePct = cats["Reduce"] || 0; + if (reducePct < 10 && reducePct > 0) { + parts.push(`TP allreduce well-optimized (${reducePct.toFixed(0)}%)`); + } + + const gemmPct = cats["GEMM"] || 0; + if (gemmPct > 20) { + parts.push(`quantize GEMMs to FP8 for ~${(gemmPct * 0.4).toFixed(0)}% improvement`); + } + + const elementPct = cats["ElementWise"] || 0; + if (elementPct > 10) { + parts.push(`fuse element-wise ops to save ~${(elementPct * 0.3).toFixed(0)}%`); + } + + if (!cudaGraph) { + parts.push("enable CUDA Graph for 3-5x speedup"); + } + + return parts.join(". ") + "."; +} + /* ── Key Findings auto-summary ── */ function KeyFindings({ kernels, kt, cudaGraph }) { diff --git a/metainfer/tasks/sglang_trace_analyze/static/sa.css b/metainfer/tasks/sglang_trace_analyze/static/sa.css index d6c47d63..097e512b 100644 --- a/metainfer/tasks/sglang_trace_analyze/static/sa.css +++ b/metainfer/tasks/sglang_trace_analyze/static/sa.css @@ -23,6 +23,11 @@ .sa-note { color: #888; font-size: 11px; margin: 4px 0; } .ml8 { margin-left: 8px; } +/* Summary banner */ +.sa-summary-banner { background: linear-gradient(135deg, #1a2a3a 0%, #1e1e1e 100%); + border: 1px solid #4a90d9; border-radius: 6px; padding: 10px 14px; + margin-bottom: 10px; font-size: 13px; color: #ddd; line-height: 1.5; } + /* Stat cards */ .sa-stats-row { display: grid; grid-template-columns: repeat(4, 1fr); gap: 10px; margin-bottom: 10px; } .sa-stat-card { background: #1e1e1e; border: 1px solid #333; border-radius: 6px;