/llm-pipeline-analysis

name: llm-pipeline-analysis
description: "Inspect LLM torch profiler traces at forward-pass, layer, and kernel level. Use when you need layer timings, anchor-kernel boundaries, representative kernel flows, or Perfetto time ranges."

LLM Pipeline Analysis

Overview

Use this when a whole-trace profiler summary is too coarse. The scripts read a Chrome-trace JSON file, find layer-boundary anchor kernels, group kernels into forward passes and layers, and print timing tables you can use for Perfetto navigation or detailed timing analysis.

When To Use It

• when you need to know **which layers** contribute most
• when the model has alternating layer types (e.g. models with

`compress_ratios` like DeepSeek-V4 NSA, or hybrid GDN/GQA stacks such as Qwen3.8-27B)

• when you need to compare cold-start vs steady-state forward passes
• when you need to navigate to a specific layer in Perfetto UI
• when you need to select representative layers for deep-dive analysis

Verify inputs from artifacts

Resolve model config, phase, rank/device and TP/EP/DP from the run manifest, server arguments and trace. Ask only for missing information that changes the analysis. `TP-0` identifies a rank, not TP world size; do not default an unknown run to TP8/EP8 or decode BS1. Speculative request BS and target-verify M differ. If no config is available, use an explicit profile and verified layer count.

Model Profiles

Scripts use **ModelProfile** to determine layer boundary detection and kernel classification. Profiles are auto-inferred from `config.json` or selected via `--profile`:

| Profile | Anchor kernel | Blocks/layer | Auto-infer condition | |---|---|---|---| | `dsv4_csa_hca` | `mhc_pre_big_fuse` | 2 | DeepSeek-V4 or compression fallback | | `dsv41` | explicit verified anchor | 1 (override for sublayers) | V4.1 text/wrapper identity or KV-source layers | | `dsv3_mla` | explicit backend anchor or residual boundary | backend-dependent | homogeneous MLA after identity dispatch | | `mla_dsa` | verified residual boundary | 2 | GLM-5/5.2/5.3, V3.2, DSA MLA | | `kimi_k3` | explicit `attn_res_fused` (SGLang) / `attn_res_fwd_online_v2_kernel` (vLLM) | 2 + one output aggregation/pass | K3 / Kimi linear with attention residual | | `glm5_next` | `mhc_pre_big_fuse` | 2 | GLM-5.3-Flash | | `hy_v4` | explicit `ihc_pre_stage2` on Triton fallback | 2 | Hy4 (hpc symbol requires trace check) | | `qwen4_exp` | explicit `grouped_gemma_rmsnorm` | 2 | Qwen3.8-Flash-Next (vLLM may fuse norm) | | `nemotron_h` | explicit residual norm | 1 | one Mamba/attention/MoE block per config layer | | `longcat_flash` | explicit residual boundary | 4 | LongCat-Flash architecture | | `gemma4` | explicit `_gemma_qkv_rmsnorm_kernel` (SGLang CUDA) | 1 | Gemma4 text/wrapper identity | | `generic_addnorm` | verified fused-add RMSNorm or allreduce fusion | 2 | other configured models | | `generic` | explicit anchor; residual candidates only | 1, or 2 for residual candidates | config-less fallback |

Nested `text_config`, `language_config` and `llm_config` are flattened before inference. Model identity takes priority over `kv_lora_rank`: K3, GLM-5 Next, Hy4, dots3 and LongCat must not be treated as homogeneous V3 MLA. Read layer count from config: V4 Flash has 43 layers, Pro 61, V4.1 40. These anchor families are source-verified for SGLang/vLLM, with backend and phase caveats; TensorRT-LLM/TokenSpeed cadence requires separate trace evidence.

Prerequisites

• A `torch.profiler` trace in Chrome-trace JSON format (`.json` or `.json.gz`)
• The model's `config.json` (for profile inference, `compress_ratios`, etc.)
• The trace must contain a recognizable layer-boundary anchor kernel

(auto-detected from the profile, or specified via `--anchor-kernel`)

Layer Boundary Detection

The scripts use an anchor kernel as a layer-boundary marker. The anchor and layer structure are determined by the active **ModelProfile**.

Find the model's residual kind at its stage boundary first. SGLang declares attention/FFN boundaries through `python/sglang/srt/layers/layer_boundary/factories.py` (add-norm, gated, iHC, mHC, post-norm and stream). Backend dispatch determines the visible boundary kernel.

V4 currently starts each half at `mhc_pre_big_fuse*`: attention-pre then FFN-pre, two anchors per layer. Fused post+pre removes `mhc_post_tilelang` for small-token SGLang traces; retain that old anchor only for verified unfused traces or the >32-token path. V4.1 uses previous-sublayer pre-mix and disables that cross-layer fused path, so it always requires explicit anchor evidence. FlashMLA's `flash_fwd_mla_combine` depends on backend and split count; SM100 SGLang V3 defaults to TRT-LLM MLA, while DSA models use DSA.

All scripts accept `--blocks-per-layer K --anchor-offset I --pid PID --device D`. `I` is a verified first-layer anchor in the selected GPU, not inferred from counts. Mixed ranks/devices and a count residue modulo `N*K` are rejected. A divisible count is necessary but insufficient: separate target/draft phases, MTP/DSPARK passes and TBO microbatches first. K3's extra output aggregation requires a phase-filtered trace for these timing analyzers; the layer-track helper supports the extra anchor explicitly. The final selected pass ends at the selected GPU event boundary, which must itself be verified as a complete pass. Anchors are navigation intervals; cross-stream kernels may overlap them.

For compiled vLLM traces, norms can lower to Inductor kernels. Capture a short mapping trace with explicit norm custom ops or eager execution and verify an attention/backend anchor on the actual measured trace.

Source audit 2026-10-05: [SGLang V4 dispatch](https://github.com/sgl-project/sglang/blob/b1bbd74f287f13ed1276b0403a01ebb55c597e93/python/sglang/srt/models/deepseek_v4.py), [mHC kernels](https://github.com/sgl-project/sglang/blob/b1bbd74f287f13ed1276b0403a01ebb55c597e93/python/sglang/kernels/ops/layernorm/mhc.py), [vLLM mHC dispatch](https://github