[Perf] Profile Kimi SM120 FlashInfer MXFP4 MoE

experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/README.md

@@ -399,3 +399,178 @@ SGLang 集成测试：`6 passed, 18 warnings in 563.41s`。warnings 为 asyncio
 2. 没有修改 checkpoint，没有 runtime monkeypatch，没有通过放宽 finite tolerance 获得通过。
 3. microbenchmark latency 仅用于检查 candidate 可执行，不作为生产性能结论。
 4. 下一阶段必须用 Nsight Systems/Compute 和三轮交错 A/B 区分 wrapper、同步、权重搬运、dispatch 与 kernel 本体；在此之前不宣称优于 Marlin。
+
+## Phase 4 算子性能与 Profiler 归因
+
+状态：**完成。** 本阶段只证明单卡本地 MoE runner 的正确性、性能边界和
+kernel 归因，不代表完整 Kimi-K3 serving 性能。真实 Scheduler、Chunked
+Prefill、跨节点 EP 通信和 CUDA Graph 的联合效果留到 Phase 5。
+
+### 实验边界
+
+两后端均通过 SGLang 生产 runner 调用，而不是直接调用底层 kernel：
+
+- baseline：`fused_experts_none_to_marlin`；
+- candidate：`fused_experts_none_to_flashinfer_mxfp4`；
+- 相同 BF16 输入、packed MXFP4 权重、E8M0 scale、Top-K IDs/weights、
+  SiTU `(alpha=4, beta=25)`；
+- Kimi 真尺寸 `H=3584, N=3072, topK=16`；
+- `EP32/rank7`，每个 rank 驻留 28 个本地专家，总专家数 896；
+- `ep_mixed` 路由中每个 token 的 8 个 slot 指向本地专家、8 个指向远端专家；
+- 仅使用 601 的物理 GPU6，并用 UUID
+  `GPU-45354dff-e8b3-1c0d-99c9-d3562bcab5fd` 做运行时绑定校验。
+
+这里的 `M` 是单次本地 MoE runner 收到的 token 数，不等于 HTTP 层的
+`ISL * concurrency`。本地平均路由项约为：
+
+```text
+local assignments / expert = M * 8 / 28
+```
+
+因此本阶段不能替代 `ISL=8192, C=16` 的真实服务测试。
+
+### 统一入口
+
+```bash
+python3 profile_kimi_sm120_mxfp4.py \
+  --shapes prefill prefill_m512 prefill_m2048 prefill_m8192 \
+  --rounds 5 --iterations 30 --warmup 5 \
+  --output results/phase4_mscan_ep32_5rounds_gpu6.json
+```
+
+正式运行时使用镜像 `lmsysorg/sglang:kimi-k3-fiv617situ-warm`，挂载：
+
+```text
+/tmp/hzy-fi-situ-phase3-v3             -> patched FlashInfer Python/AOT package
+/data/hzy/src/sglang-kimi-sm120        -> patched SGLang source
+/data/hzy/cache/sglang-jit-sm120       -> persistent SGLang JIT cache
+```
+
+并设置：
+
+```text
+FLASHINFER_DISABLE_JIT=1
+FLASHINFER_DISABLE_VERSION_CHECK=1
+SGLANG_JIT_CACHE_DIR=/root/.cache/sglang/jit
+EXPECTED_GPU_UUID=45354dff-e8b3-1c0d-99c9-d3562bcab5fd
+```
+
+### 五轮交错 A/B
+
+每个 shape 先做 BF16 参考正确性，再 warmup 5 次。五轮按
+`Marlin -> FlashInfer`、`FlashInfer -> Marlin` 交替执行，每轮每后端 30 次。
+下表为 GPU latency 中位数：
+
+| M | 平均本地路由项/专家 | Marlin | FlashInfer | FlashInfer 相对速度 | FI vs reference |
+|---:|---:|---:|---:|---:|---:|
+| 128 | 36.6 | 0.870 ms | 0.926 ms | 0.94x，慢约 6.5% | cosine 0.999615，NRMSE 2.774% |
+| 512 | 146.3 | 2.610 ms | 1.782 ms | 1.46x | cosine 0.999614，NRMSE 2.779% |
+| 2048 | 585.1 | 9.070 ms | 4.355 ms | 2.08x | cosine 0.999623，NRMSE 2.746% |
+| 8192 | 2340.6 | 34.379 ms | 16.167 ms | 2.13x | cosine 0.999623，NRMSE 2.744% |
+
+五轮逐轮比值稳定：M512 为 `1.428x-1.474x`，M2048 为
+`2.065x-2.089x`，M8192 为 `2.126x-2.128x`。这证明 FlashInfer 的优势
+依赖每次 runner 的实际 M；它不是小 batch 的无条件替代品。
+
+M512 的第一次 pilot 只 warmup 1 次，得到 FlashInfer 4.26 ms、Marlin
+1.94 ms，与正式结果相反。该结果保留为 warmup 敏感性证据，不用于性能结论。
+
+### Nsight Systems
+
+捕获文件：
+
+```text
+results/phase4_nsys_ab_gpu6.nsys-rep
+results/phase4_nsys_ab_gpu6.log
+results/phase4_nsys_stats_*.csv
+results/phase4_nsys_*_iteration_5_*.csv
+```
+
+稳定的第 5 次迭代中，主 GEMM kernel 总时间：
+
+| Shape | Marlin | FlashInfer | 观察 |
+|---|---:|---:|---|
+| Decode M1 | 109.4 us | 191.8 us | FlashInfer 主 kernel 约慢 75% |
+| Prefill M128 | 709.9 us | 865.3 us | FlashInfer 主 kernel 约慢 22% |
+
+Decode 的非主 kernel 开销约为 FlashInfer 20.3 us、Marlin 6.4 us；Prefill
+分别约 45.8 us、44.2 us。小 M 差距主要来自 kernel 本体，不能归因成纯
+Python wrapper 或同步开销。
+
+### Nsight Compute
+
+RTX 6000D 需要较新的工具。镜像内 NCU 2025.3.1 返回
+`Profiling is not supported on device 0`；正式报告由 NVIDIA 官方包
+Nsight Compute 2026.2.1 生成：
+
+```text
+/data/hzy/tools/nsight-compute-2026.2.1/
+results/phase4_ncu_2026_2_1_ab_gpu6.ncu-rep
+results/phase4_ncu_2026_2_1_ab_gpu6.log
+results/phase4_ncu_2026_2_1_ab_gpu6_raw.csv
+```
+
+`--set basic` 捕获到的两组主 kernel 指标：
+
+| Shape/backend | 主 kernel 合计 | DRAM 峰值利用率 | SM 峰值利用率 | registers/thread | shared memory/block |
+|---|---:|---:|---:|---:|---:|
+| Decode Marlin | 121.2 us | 77.8%-86.8% | 21.2%-22.5% | 113 | 33.8 KiB |
+| Decode FlashInfer | 199.4 us | 49.9%-51.7% | 47.7%-47.8% | 168 | 84.0 KiB |
+| M128 Marlin | 513.0 us | 73.7%-76.1% | 61.1%-65.9% | 255 | 102.4 KiB |
+| M128 FlashInfer | 535.1 us | 67.2%-70.7% | 61.1%-63.3% | 168 | 84.0 KiB |
+
+NCU 使用 replay，绝对时间不与 Nsight Systems 混用。它支持相同方向的
+结论：Decode 时 FlashInfer 使用更多寄存器和共享内存，DRAM 利用率更低，
+当前 CUTLASS 配置没有在小 M 上胜过 Marlin；M128 的主 kernel 已接近，
+但完整 runner 仍有其他步骤。
+
+### 工程问题与修复记录
+
+1. `--privileged` 会重新暴露全部 `/dev/nvidia*`，曾令容器落到物理 GPU0。
+   UUID guard 立即终止实验。之后固定使用 `--gpus device=6`、
+   `NVIDIA_VISIBLE_DEVICES=6`，不再使用 privileged。
+2. 纯 Python wheel 在每个新进程触发 FlashInfer JIT；NCU 注入子进程后会
+   卡住大量 `nvcc`。正式 profiling 把三个已验证 `.so` 放入
+   `flashinfer/data/aot/`，并设置 `FLASHINFER_DISABLE_JIT=1`。
+3. SGLang 自身 JIT cache 持久化到 `/data/hzy/cache/sglang-jit-sm120`。
+4. `torch.quantile` 对 M8192 的约 2936 万元素输出报输入过大。统计代码改用
+   精确 `torch.kthvalue`；cosine、NRMSE、max/mean 与正确性阈值均未放宽。
+
+详细时间线和失败证据见 `decision_log.md`。
+
+### Phase 4 决策
+
+1. FlashInfer SiTU MXFP4 路径已通过正确性门槛，并在 M512 及以上的本地
+   runner shape 显示明确吞吐潜力。
+2. 它在 Decode/M128 小 M 上仍落后于 Marlin，不能现在就设为无条件默认。
+3. 不再扩展 synthetic EP4/kernel sweep。下一步应进入四节点真实 Kimi-K3
+   serving，观察 Scheduler 实际 M 分布及端到端 TTFT/TPOT/TPS。
+4. 在 Phase 5 完成前，不宣称 FlashInfer 提升生产性能；当前 PR 的可靠表述
+   是补齐 Kimi-K3 + SM120 + SiTU + MXFP4 功能支持，并提供显式 Marlin fallback。
+
+### Phase 4 原始证据
+
+服务器根目录：
+
+```text
+/data/hzy/sskj/experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/
+```
+
+关键文件：
+
+```text
+phase4_ab_5rounds_gpu6.json
+phase4_mscan_ep32_5rounds_gpu6.json
+phase4_mscan_ep32_5rounds_gpu6.log
+phase4_aot_nojit_smoke_gpu6.json
+phase4_nsys_ab_gpu6.nsys-rep
+phase4_nsys_stats_*.csv
+phase4_ncu_2026_2_1_ab_gpu6.ncu-rep
+phase4_ncu_2026_2_1_ab_gpu6_raw.csv
+```
+
+本地备份：
+
+```text
+/Users/hzy/Desktop/infra/.phase3_work/phase4_results/
+```

experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/decision_log.md

@@ -0,0 +1,59 @@
+# Kimi-K3 SM120 FlashInfer MXFP4 Decision Log
+
+本文件记录可恢复的关键决策、失败尝试和证据路径。成功结论见 `README.md`。
+
+## 2026-08-18 Phase 4
+
+### 对称 A/B
+
+- 决策：baseline/candidate 都走 SGLang production runner，共享输入、权重、
+  scale、routing、SiTU 参数和计时方式。
+- 结果：M1/M128 FlashInfer 略慢；扩展五轮扫描中 M512/2048/8192 分别比
+  Marlin 快约 1.46x/2.08x/2.13x。
+- 限制：M 是单卡本地 runner token 数，不是 HTTP ISL 或全局 batch。
+- 证据：`results/phase4_ab_5rounds_gpu6.json`、
+  `results/phase4_mscan_ep32_5rounds_gpu6.json`。
+
+### GPU 绑定事故
+
+- 现象：最初 Nsight 命令同时使用 `--gpus device=6` 和 `--privileged`，UUID
+  guard 检测到实际设备为物理 GPU0 并终止。
+- 根因：privileged 重新暴露全部 NVIDIA device node，破坏容器 GPU 隔离。
+- 处理：移除 privileged；固定 `--gpus device=6`、
+  `NVIDIA_VISIBLE_DEVICES=6` 和预期 UUID。未继续占用 GPU0。
+
+### JIT 与 AOT
+
+- 现象：纯 Python FlashInfer wheel 在新进程重新 JIT。NCU profiler 注入
+  JIT 子进程后产生大量挂起 `nvcc`。
+- 处理：将已验证的 `fused_moe_120`、`mxfp8_quantization_sm100`、
+  `fp4_quantization_120f` `.so` 写入 wheel 的 AOT 目录；设置
+  `FLASHINFER_DISABLE_JIT=1`；持久化 SGLang JIT cache。
+- 验证：`results/phase4_aot_nojit_smoke_gpu6.json`。
+
+### Profiler 工具
+
+- NCU 2025.3.1：RTX 6000D/SM120 不受支持，报
+  `Profiling is not supported on device 0`。
+- NCU 2026.2.1：从 NVIDIA 官方 CUDA repository 下载并校验 SHA256，解压到
+  `/data/hzy/tools/nsight-compute-2026.2.1/`，正式捕获成功。
+- Nsight Systems：正式捕获成功，原始文件为
+  `results/phase4_nsys_ab_gpu6.nsys-rep`。
+
+### 大输出统计
+
+- 现象：M8192 输出约 2936 万元素，`torch.quantile` 报 input too large。
+- 处理：P95/P99 改用精确 `torch.kthvalue`。未抽样，未改正确性阈值。
+
+### Warmup 敏感性
+
+- M512 pilot 使用 warmup=1，得到与正式结果相反的排序。
+- 正式五轮使用 warmup=5、每轮 30 次，五轮速度比为
+  `1.428x-1.474x`，采用正式结果。
+- pilot 日志保留，不作为性能结论。
+
+### 阶段出口
+
+- 不继续 synthetic EP4 扫描。
+- 下一步需要 601-604 共 32 卡运行真实 TP32/EP32 Kimi-K3 serving。
+- 未经资源授权不清理现有作业、不启动 Phase 5。

experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/profile_kimi_sm120_mxfp4.py

@@ -0,0 +1,583 @@
+#!/usr/bin/env python3
+"""Profile Kimi-K3 MXFP4 MoE on SM120: SGLang Marlin vs FlashInfer.
+
+Both backends consume the same BF16 input, packed MXFP4 checkpoint values,
+E8M0 scales, global top-k ids/weights, SiTU parameters, and EP32/rank7 mapping.
+Backend-specific post-load layouts are derived once before correctness checks,
+warmup, timing, and profiling.
+"""
+
+from __future__ import annotations
+
+import argparse
+from contextlib import nullcontext
+from dataclasses import asdict
+from datetime import datetime
+import json
+import math
+import os
+from pathlib import Path
+import statistics
+import time
+from types import SimpleNamespace
+from typing import Callable
+
+import torch
+import torch.nn.functional as F
+
+from flashinfer import (
+    SfLayout,
+    block_scale_interleave,
+    mxfp4_dequantize,
+    mxfp4_quantize,
+    mxfp8_quantize,
+)
+from sglang.srt.distributed import init_distributed_environment
+from sglang.srt.distributed.parallel_state import (
+    destroy_model_parallel,
+    initialize_model_parallel,
+)
+from sglang.srt.layers.moe.moe_runner.base import MoeRunnerConfig
+from sglang.srt.layers.moe.moe_runner.flashinfer_cutlass import (
+    FlashInferCutlassMxfp4MoeQuantInfo,
+    fused_experts_none_to_flashinfer_mxfp4,
+)
+from sglang.srt.layers.moe.moe_runner.marlin import (
+    MarlinMoeQuantInfo,
+    fused_experts_none_to_marlin,
+)
+from sglang.srt.layers.moe.token_dispatcher.standard import StandardDispatchOutput
+from sglang.srt.layers.moe.topk import StandardTopKOutput
+from sglang.srt.layers.quantization.marlin_utils_fp4 import (
+    prepare_moe_mxfp4_layer_for_marlin,
+)
+
+from validate_kimi_sm120_mxfp4 import (
+    Case,
+    QuantizedCase,
+    dequantize_mxfp8_input,
+    make_routing,
+    run_reference,
+)
+
+
+SHAPES = {
+    "decode": Case(
+        "kimi_decode_m1_ep32",
+        1,
+        3584,
+        3072,
+        28,
+        16,
+        "ep_mixed",
+        ep_size=32,
+        ep_rank=7,
+    ),
+    "prefill": Case(
+        "kimi_prefill_m128_ep32",
+        128,
+        3584,
+        3072,
+        28,
+        16,
+        "ep_mixed",
+        ep_size=32,
+        ep_rank=7,
+    ),
+    "prefill_m512": Case(
+        "kimi_prefill_m512_ep32",
+        512,
+        3584,
+        3072,
+        28,
+        16,
+        "ep_mixed",
+        ep_size=32,
+        ep_rank=7,
+    ),
+    "prefill_m2048": Case(
+        "kimi_prefill_m2048_ep32",
+        2048,
+        3584,
+        3072,
+        28,
+        16,
+        "ep_mixed",
+        ep_size=32,
+        ep_rank=7,
+    ),
+    "prefill_m8192": Case(
+        "kimi_prefill_m8192_ep32",
+        8192,
+        3584,
+        3072,
+        28,
+        16,
+        "ep_mixed",
+        ep_size=32,
+        ep_rank=7,
+    ),
+}
+
+
+def _quantize_experts_linear(
+    weight: torch.Tensor,
+) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
+    packed = []
+    scales = []
+    dequantized = []
+    for expert_weight in weight:
+        expert_packed, expert_scale = mxfp4_quantize(
+            expert_weight,
+            sfLayout=SfLayout.layout_linear,
+        )
+        packed.append(expert_packed)
+        scales.append(expert_scale)
+        dequantized.append(
+            mxfp4_dequantize(
+                expert_packed.cpu(),
+                expert_scale.cpu(),
+                sfLayout=SfLayout.layout_linear,
+            )
+        )
+    return (
+        torch.stack(packed),
+        torch.stack(scales),
+        torch.stack(dequantized).to(device=weight.device, dtype=weight.dtype),
+    )
+
+
+def _prepare_linear_checkpoint_case(
+    case: Case,
+    seed: int,
+    device: torch.device,
+) -> QuantizedCase:
+    generator = torch.Generator(device=device).manual_seed(seed)
+    dtype = torch.bfloat16
+    x = torch.randn(
+        case.tokens,
+        case.hidden,
+        generator=generator,
+        dtype=dtype,
+        device=device,
+    )
+    gate = torch.randn(
+        case.local_experts,
+        case.intermediate,
+        case.hidden,
+        generator=generator,
+        dtype=dtype,
+        device=device,
+    ) * (1.0 / math.sqrt(case.hidden))
+    up = torch.randn(
+        case.local_experts,
+        case.intermediate,
+        case.hidden,
+        generator=generator,
+        dtype=dtype,
+        device=device,
+    ) * (1.0 / math.sqrt(case.hidden))
+    down = torch.randn(
+        case.local_experts,
+        case.hidden,
+        case.intermediate,
+        generator=generator,
+        dtype=dtype,
+        device=device,
+    ) * (1.0 / math.sqrt(case.intermediate))
+    input_q, input_sf = mxfp8_quantize(
+        x,
+        is_sf_swizzled_layout=True,
+        alignment=32,
+    )
+    input_dq = dequantize_mxfp8_input(input_q, input_sf, dtype)
+    fc1_q, fc1_sf, fc1_dq = _quantize_experts_linear(
+        torch.cat((up, gate), dim=1).contiguous()
+    )
+    fc2_q, fc2_sf, fc2_dq = _quantize_experts_linear(down)
+    topk_ids, topk_weights = make_routing(case, device)
+    return QuantizedCase(
+        case=case,
+        input_q=input_q,
+        input_sf=input_sf,
+        input_dq=input_dq,
+        fc1_q=fc1_q,
+        fc1_sf=fc1_sf,
+        fc1_dq=fc1_dq,
+        fc2_q=fc2_q,
+        fc2_sf=fc2_sf,
+        fc2_dq=fc2_dq,
+        topk_ids=topk_ids,
+        topk_weights=topk_weights,
+        expert_scale=torch.ones(
+            case.local_experts, dtype=torch.float32, device=device
+        ),
+        situ_gate_beta=torch.full(
+            (case.local_experts,), 4.0, dtype=torch.float32, device=device
+        ),
+        situ_linear_beta=torch.full(
+            (case.local_experts,), 25.0, dtype=torch.float32, device=device
+        ),
+    )
+
+
+def _init_sglang_distributed() -> None:
+    """Create the rank-one groups that production MoE wrappers expect."""
+    torch.cuda.set_device(0)
+    init_distributed_environment(
+        world_size=1,
+        rank=0,
+        local_rank=0,
+        distributed_init_method="tcp://127.0.0.1:29541",
+        backend="nccl",
+    )
+    initialize_model_parallel(
+        tensor_model_parallel_size=1,
+        expert_model_parallel_size=1,
+    )
+
+
+def _destroy_sglang_distributed() -> None:
+    destroy_model_parallel()
+    if torch.distributed.is_initialized():
+        torch.distributed.destroy_process_group()
+
+
+def _quantile_exact(values: torch.Tensor, quantile: float) -> float:
+    """Return an exact quantile without torch.quantile's 2^24 element limit."""
+    rank = max(1, math.ceil(quantile * values.numel()))
+    return float(torch.kthvalue(values, rank).values.item())
+
+
+def _metrics(reference: torch.Tensor, candidate: torch.Tensor) -> dict:
+    ref = reference.float().reshape(-1)
+    got = candidate.float().reshape(-1)
+    diff = got - ref
+    abs_diff = diff.abs()
+    denominator = torch.sqrt(torch.mean(ref.square())).clamp_min(1e-12)
+    return {
+        "mean_abs": float(abs_diff.mean().item()),
+        "p95_abs": _quantile_exact(abs_diff, 0.95),
+        "p99_abs": _quantile_exact(abs_diff, 0.99),
+        "max_abs": float(abs_diff.max().item()),
+        "nrmse": float((torch.sqrt(torch.mean(diff.square())) / denominator).item()),
+        "cosine": float(F.cosine_similarity(ref, got, dim=0).item()),
+        "reference_rms": float(denominator.item()),
+        "reference_max_abs": float(ref.abs().max().item()),
+        "finite": bool(torch.isfinite(candidate).all()),
+    }
+
+
+def _assert_close(name: str, reference: torch.Tensor, candidate: torch.Tensor) -> dict:
+    metrics = _metrics(reference, candidate)
+    if not metrics["finite"]:
+        raise AssertionError(f"{name} produced non-finite output: {metrics}")
+    torch.testing.assert_close(candidate, reference, rtol=0.2, atol=0.2)
+    return metrics
+
+
+def _checkpoint_gate_up_from_flashinfer_up_gate(tensor: torch.Tensor) -> torch.Tensor:
+    up, gate = tensor.chunk(2, dim=1)
+    return torch.cat((gate, up), dim=1).contiguous()
+
+
+def _build_expert_map(case: Case, device: torch.device) -> torch.Tensor:
+    expert_map = torch.full(
+        (case.total_experts,), -1, dtype=torch.int32, device=device
+    )
+    start = case.local_expert_start
+    expert_map[start : start + case.local_experts] = torch.arange(
+        case.local_experts, dtype=torch.int32, device=device
+    )
+    return expert_map
+
+
+class BackendPair:
+    def __init__(self, case: Case, seed: int):
+        self.case = case
+        self.data = _prepare_linear_checkpoint_case(
+            case,
+            seed,
+            torch.device("cuda"),
+        )
+        self.config = MoeRunnerConfig(
+            num_experts=case.total_experts,
+            num_local_experts=case.local_experts,
+            hidden_size=case.hidden,
+            intermediate_size_per_partition=case.intermediate,
+            top_k=case.top_k,
+            activation="situ",
+            is_gated=True,
+            gemm1_alpha=4.0,
+            gemm1_clamp_limit=25.0,
+            gate_up_interleaved=False,
+        )
+        logits = torch.zeros(
+            case.tokens,
+            case.total_experts,
+            dtype=torch.float32,
+            device="cuda",
+        )
+        self.expert_map = _build_expert_map(case, torch.device("cuda"))
+        self.flashinfer_dispatch = StandardDispatchOutput(
+            self.data.input_dq.contiguous(),
+            None,
+            StandardTopKOutput(
+                self.data.topk_weights,
+                self.data.topk_ids,
+                logits,
+            ),
+        )
+        marlin_topk_ids = self.expert_map[self.data.topk_ids]
+        self.marlin_dispatch = StandardDispatchOutput(
+            self.data.input_dq.contiguous(),
+            None,
+            StandardTopKOutput(
+                self.data.topk_weights,
+                marlin_topk_ids,
+                logits,
+            ),
+        )
+        self.reference = run_reference(self.data, "situ")
+        self.fi_quant = self._build_flashinfer_quant_info()
+        self.marlin_quant = self._build_marlin_quant_info()
+
+    def _build_flashinfer_quant_info(self) -> FlashInferCutlassMxfp4MoeQuantInfo:
+        case = self.case
+        w13_scale = block_scale_interleave(self.data.fc1_sf).reshape_as(
+            self.data.fc1_sf
+        )
+        w2_scale = block_scale_interleave(self.data.fc2_sf).reshape_as(
+            self.data.fc2_sf
+        )
+        return FlashInferCutlassMxfp4MoeQuantInfo(
+            w13_weight=self.data.fc1_q,
+            w2_weight=self.data.fc2_q,
+            w13_weight_scale=w13_scale.contiguous(),
+            w2_weight_scale=w2_scale.contiguous(),
+            mxfp4_weight_global_scale=self.data.expert_scale,
+            w13_bias=None,
+            w2_bias=None,
+            swiglu_alpha=self.data.situ_gate_beta,
+            swiglu_beta=self.data.situ_linear_beta,
+            swiglu_limit=None,
+            moe_tp_size=1,
+            moe_tp_rank=0,
+            moe_ep_size=case.ep_size,
+            moe_ep_rank=case.ep_rank,
+            padded_hidden=case.hidden,
+        )
+
+    def _build_marlin_quant_info(self) -> MarlinMoeQuantInfo:
+        case = self.case
+        layer = SimpleNamespace(
+            w13_weight=torch.nn.Parameter(
+                _checkpoint_gate_up_from_flashinfer_up_gate(self.data.fc1_q.clone()),
+                requires_grad=False,
+            ),
+            w2_weight=torch.nn.Parameter(
+                self.data.fc2_q.clone(), requires_grad=False
+            ),
+            w13_weight_scale=torch.nn.Parameter(
+                _checkpoint_gate_up_from_flashinfer_up_gate(
+                    self.data.fc1_sf.clone()
+                ),
+                requires_grad=False,
+            ),
+            w2_weight_scale=torch.nn.Parameter(
+                self.data.fc2_sf.clone(), requires_grad=False
+            ),
+            orig_dtype=torch.bfloat16,
+        )
+        prepare_moe_mxfp4_layer_for_marlin(layer)
+        return MarlinMoeQuantInfo(
+            w13_qweight=layer.w13_weight,
+            w2_qweight=layer.w2_weight,
+            w13_scales=layer.w13_weight_scale,
+            w2_scales=layer.w2_weight_scale,
+            w13_g_idx_sort_indices=None,
+            w2_g_idx_sort_indices=None,
+            weight_bits=4,
+            is_k_full=True,
+            expert_map=self.expert_map,
+            global_num_experts=case.total_experts,
+        )
+
+    def flashinfer(self) -> torch.Tensor:
+        return fused_experts_none_to_flashinfer_mxfp4(
+            self.flashinfer_dispatch, self.fi_quant, self.config
+        ).hidden_states
+
+    def marlin(self) -> torch.Tensor:
+        return fused_experts_none_to_marlin(
+            self.marlin_dispatch, self.marlin_quant, self.config
+        ).hidden_states[:, : self.case.hidden]
+
+    def validate(self) -> dict:
+        fi = self.flashinfer()
+        marlin = self.marlin()
+        torch.cuda.synchronize()
+        metrics = {
+            "flashinfer_vs_reference": _metrics(self.reference, fi),
+            "marlin_vs_reference": _metrics(self.reference, marlin),
+            "flashinfer_vs_marlin": _metrics(marlin, fi),
+        }
+        print("CORRECTNESS " + json.dumps(metrics, sort_keys=True), flush=True)
+        _assert_close("flashinfer", self.reference, fi)
+        _assert_close("marlin", self.reference, marlin)
+        _assert_close("cross_backend", marlin, fi)
+        return metrics
+
+
+def _measure(fn: Callable[[], torch.Tensor], iterations: int) -> dict:
+    start = torch.cuda.Event(enable_timing=True)
+    end = torch.cuda.Event(enable_timing=True)
+    torch.cuda.synchronize()
+    wall_start = time.perf_counter()
+    start.record()
+    for _ in range(iterations):
+        fn()
+    end.record()
+    torch.cuda.synchronize()
+    wall_ms = (time.perf_counter() - wall_start) * 1000.0 / iterations
+    return {
+        "gpu_ms": start.elapsed_time(end) / iterations,
+        "wall_ms": wall_ms,
+    }
+
+
+def _warmup(pair: BackendPair, iterations: int) -> None:
+    for _ in range(iterations):
+        pair.marlin()
+        pair.flashinfer()
+    torch.cuda.synchronize()
+
+
+def run_benchmark(args: argparse.Namespace) -> dict:
+    report = {
+        "created_at": datetime.now().astimezone().isoformat(),
+        "device": torch.cuda.get_device_name(),
+        "compute_capability": list(torch.cuda.get_device_capability()),
+        "cuda_visible_devices": os.environ.get("CUDA_VISIBLE_DEVICES"),
+        "rounds": args.rounds,
+        "iterations": args.iterations,
+        "shapes": {},
+    }
+    for shape_name in args.shapes:
+        pair = BackendPair(SHAPES[shape_name], args.seed)
+        correctness = pair.validate()
+        _warmup(pair, args.warmup)
+        rounds = []
+        for round_index in range(args.rounds):
+            order = (
+                ("marlin", "flashinfer")
+                if round_index % 2 == 0
+                else ("flashinfer", "marlin")
+            )
+            result = {"round": round_index + 1, "order": list(order)}
+            for backend in order:
+                result[backend] = _measure(getattr(pair, backend), args.iterations)
+            rounds.append(result)
+        summary = {}
+        for backend in ("marlin", "flashinfer"):
+            summary[backend] = {
+                metric: statistics.median(
+                    result[backend][metric] for result in rounds
+                )
+                for metric in ("gpu_ms", "wall_ms")
+            }
+        summary["speedup_marlin_over_flashinfer"] = {
+            metric: summary["marlin"][metric] / summary["flashinfer"][metric]
+            for metric in ("gpu_ms", "wall_ms")
+        }
+        report["shapes"][shape_name] = {
+            "case": asdict(pair.case),
+            "correctness": correctness,
+            "rounds": rounds,
+            "summary": summary,
+        }
+        del pair
+        torch.cuda.empty_cache()
+    return report
+
+
+def run_profile(args: argparse.Namespace) -> None:
+    shape_names = tuple(SHAPES) if args.shape == "all" else (args.shape,)
+    pairs = [
+        (shape_name, BackendPair(SHAPES[shape_name], args.seed))
+        for shape_name in shape_names
+    ]
+    for shape_name, pair in pairs:
+        print(
+            "PROFILE_CORRECTNESS "
+            + json.dumps({shape_name: pair.validate()}, sort_keys=True),
+            flush=True,
+        )
+        _warmup(pair, args.warmup)
+    torch.cuda.synchronize()
+    torch.cuda.cudart().cudaProfilerStart()
+    backends = (
+        (args.backend,)
+        if args.backend != "both"
+        else ("marlin", "flashinfer")
+    )
+    for shape_name, pair in pairs:
+        for backend in backends:
+            for index in range(args.profile_iterations):
+                label = f"{backend}_{shape_name}_iteration_{index + 1}"
+                context = (
+                    torch.cuda.nvtx.range(label)
+                    if torch.cuda.is_available()
+                    else nullcontext()
+                )
+                with context:
+                    getattr(pair, backend)()
+    torch.cuda.synchronize()
+    torch.cuda.cudart().cudaProfilerStop()
+
+
+def parse_args() -> argparse.Namespace:
+    parser = argparse.ArgumentParser()
+    parser.add_argument("--mode", choices=("benchmark", "profile"), default="benchmark")
+    parser.add_argument("--shapes", nargs="+", choices=tuple(SHAPES), default=list(SHAPES))
+    parser.add_argument("--shape", choices=(*SHAPES, "all"), default="decode")
+    parser.add_argument("--backend", choices=("marlin", "flashinfer", "both"), default="both")
+    parser.add_argument("--rounds", type=int, default=3)
+    parser.add_argument("--iterations", type=int, default=30)
+    parser.add_argument("--warmup", type=int, default=5)
+    parser.add_argument("--profile-iterations", type=int, default=3)
+    parser.add_argument("--seed", type=int, default=20260817)
+    parser.add_argument("--output", type=Path)
+    return parser.parse_args()
+
+
+def main() -> None:
+    args = parse_args()
+    if not torch.cuda.is_available():
+        raise SystemExit("CUDA is required")
+    expected_uuid = os.environ.get("EXPECTED_GPU_UUID")
+    actual_uuid = str(torch.cuda.get_device_properties(0).uuid)
+    if expected_uuid and actual_uuid.lower() != expected_uuid.lower():
+        raise SystemExit(
+            f"GPU binding mismatch: expected {expected_uuid}, got {actual_uuid}"
+        )
+    if torch.cuda.get_device_capability() != (12, 0):
+        raise SystemExit(f"SM120 is required, got {torch.cuda.get_device_capability()}")
+    _init_sglang_distributed()
+    try:
+        with torch.inference_mode():
+            if args.mode == "profile":
+                run_profile(args)
+                return
+            report = run_benchmark(args)
+    finally:
+        _destroy_sglang_distributed()
+    rendered = json.dumps(report, indent=2, sort_keys=True)
+    print(rendered)
+    if args.output is not None:
+        args.output.parent.mkdir(parents=True, exist_ok=True)
+        args.output.write_text(rendered + "\n", encoding="utf-8")
+
+
+if __name__ == "__main__":
+    main()

experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/phase4_ab_5rounds_gpu6.json

@@ -0,0 +1,293 @@
+{
+  "compute_capability": [
+    12,
+    0
+  ],
+  "created_at": "2026-08-17T09:47:32.426483+00:00",
+  "cuda_visible_devices": null,
+  "device": "NVIDIA RTX 6000D",
+  "iterations": 100,
+  "rounds": 5,
+  "shapes": {
+    "decode": {
+      "case": {
+        "ep_rank": 7,
+        "ep_size": 32,
+        "hidden": 3584,
+        "inject_nonfinite": false,
+        "intermediate": 3072,
+        "local_experts": 28,
+        "name": "kimi_decode_m1_ep32",
+        "routing": "ep_mixed",
+        "tokens": 1,
+        "top_k": 16
+      },
+      "correctness": {
+        "flashinfer_vs_marlin": {
+          "cosine": 0.9996320605278015,
+          "finite": true,
+          "max_abs": 0.00978851318359375,
+          "mean_abs": 0.002149919280782342,
+          "nrmse": 0.02712450921535492,
+          "p95_abs": 0.00537109375,
+          "p99_abs": 0.007080078125,
+          "reference_max_abs": 0.37890625,
+          "reference_rms": 0.09998540580272675
+        },
+        "flashinfer_vs_reference": {
+          "cosine": 0.9996223449707031,
+          "finite": true,
+          "max_abs": 0.009765625,
+          "mean_abs": 0.00218405993655324,
+          "nrmse": 0.027483593672513962,
+          "p95_abs": 0.00537109375,
+          "p99_abs": 0.0072021484375,
+          "reference_max_abs": 0.37890625,
+          "reference_rms": 0.10006484389305115
+        },
+        "marlin_vs_reference": {
+          "cosine": 0.9999805688858032,
+          "finite": true,
+          "max_abs": 0.00390625,
+          "mean_abs": 0.00045798509381711483,
+          "nrmse": 0.0062673031352460384,
+          "p95_abs": 0.0009765625,
+          "p99_abs": 0.001953125,
+          "reference_max_abs": 0.37890625,
+          "reference_rms": 0.10006484389305115
+        }
+      },
+      "rounds": [
+        {
+          "flashinfer": {
+            "gpu_ms": 0.3265369415283203,
+            "wall_ms": 0.32803289126604795
+          },
+          "marlin": {
+            "gpu_ms": 0.30863103866577146,
+            "wall_ms": 0.3124843700788915
+          },
+          "order": [
+            "marlin",
+            "flashinfer"
+          ],
+          "round": 1
+        },
+        {
+          "flashinfer": {
+            "gpu_ms": 0.3199043273925781,
+            "wall_ms": 0.32471023965626955
+          },
+          "marlin": {
+            "gpu_ms": 0.3098236846923828,
+            "wall_ms": 0.3145282808691263
+          },
+          "order": [
+            "flashinfer",
+            "marlin"
+          ],
+          "round": 2
+        },
+        {
+          "flashinfer": {
+            "gpu_ms": 0.32463966369628905,
+            "wall_ms": 0.3285454004071653
+          },
+          "marlin": {
+            "gpu_ms": 0.30631423950195313,
+            "wall_ms": 0.3102407115511596
+          },
+          "order": [
+            "marlin",
+            "flashinfer"
+          ],
+          "round": 3
+        },
+        {
+          "flashinfer": {
+            "gpu_ms": 0.31916448593139646,
+            "wall_ms": 0.32354930881410837
+          },
+          "marlin": {
+            "gpu_ms": 0.3083350372314453,
+            "wall_ms": 0.3130973596125841
+          },
+          "order": [
+            "flashinfer",
+            "marlin"
+          ],
+          "round": 4
+        },
+        {
+          "flashinfer": {
+            "gpu_ms": 0.3248342514038086,
+            "wall_ms": 0.326665451284498
+          },
+          "marlin": {
+            "gpu_ms": 0.31008224487304686,
+            "wall_ms": 0.31198404962196946
+          },
+          "order": [
+            "marlin",
+            "flashinfer"
+          ],
+          "round": 5
+        }
+      ],
+      "summary": {
+        "flashinfer": {
+          "gpu_ms": 0.32463966369628905,
+          "wall_ms": 0.326665451284498
+        },
+        "marlin": {
+          "gpu_ms": 0.30863103866577146,
+          "wall_ms": 0.3124843700788915
+        },
+        "speedup_marlin_over_flashinfer": {
+          "gpu_ms": 0.9506880186843276,
+          "wall_ms": 0.9565883654061232
+        }
+      }
+    },
+    "prefill": {
+      "case": {
+        "ep_rank": 7,
+        "ep_size": 32,
+        "hidden": 3584,
+        "inject_nonfinite": false,
+        "intermediate": 3072,
+        "local_experts": 28,
+        "name": "kimi_prefill_m128_ep32",
+        "routing": "ep_mixed",
+        "tokens": 128,
+        "top_k": 16
+      },
+      "correctness": {
+        "flashinfer_vs_marlin": {
+          "cosine": 0.9996392726898193,
+          "finite": true,
+          "max_abs": 0.0126953125,
+          "mean_abs": 0.0021022639703005552,
+          "nrmse": 0.02686966024339199,
+          "p95_abs": 0.005126953125,
+          "p99_abs": 0.0068359375,
+          "reference_max_abs": 0.462890625,
+          "reference_rms": 0.09850681573152542
+        },
+        "flashinfer_vs_reference": {
+          "cosine": 0.9996151924133301,
+          "finite": true,
+          "max_abs": 0.013671875,
+          "mean_abs": 0.00217337254434824,
+          "nrmse": 0.02774309180676937,
+          "p95_abs": 0.00537109375,
+          "p99_abs": 0.00689697265625,
+          "reference_max_abs": 0.46484375,
+          "reference_rms": 0.09856347739696503
+        },
+        "marlin_vs_reference": {
+          "cosine": 0.9999709129333496,
+          "finite": true,
+          "max_abs": 0.00390625,
+          "mean_abs": 0.000566127069760114,
+          "nrmse": 0.007645560894161463,
+          "p95_abs": 0.00146484375,
+          "p99_abs": 0.001953125,
+          "reference_max_abs": 0.46484375,
+          "reference_rms": 0.09856347739696503
+        }
+      },
+      "rounds": [
+        {
+          "flashinfer": {
+            "gpu_ms": 0.9432563018798829,
+            "wall_ms": 0.9479056694544852
+          },
+          "marlin": {
+            "gpu_ms": 0.8844432067871094,
+            "wall_ms": 0.8977009798400104
+          },
+          "order": [
+            "marlin",
+            "flashinfer"
+          ],
+          "round": 1
+        },
+        {
+          "flashinfer": {
+            "gpu_ms": 0.9416944122314453,
+            "wall_ms": 0.9463393897749484
+          },
+          "marlin": {
+            "gpu_ms": 0.88412353515625,
+            "wall_ms": 0.8888133312575519
+          },
+          "order": [
+            "flashinfer",
+            "marlin"
+          ],
+          "round": 2
+        },
+        {
+          "flashinfer": {
+            "gpu_ms": 0.9426467132568359,
+            "wall_ms": 0.9474158892408013
+          },
+          "marlin": {
+            "gpu_ms": 0.8948384094238281,
+            "wall_ms": 0.8993784594349563
+          },
+          "order": [
+            "marlin",
+            "flashinfer"
+          ],
+          "round": 3
+        },
+        {
+          "flashinfer": {
+            "gpu_ms": 0.9381990051269531,
+            "wall_ms": 0.9428730909712613
+          },
+          "marlin": {
+            "gpu_ms": 0.89110595703125,
+            "wall_ms": 0.8956197183579206
+          },
+          "order": [
+            "flashinfer",
+            "marlin"
+          ],
+          "round": 4
+        },
+        {
+          "flashinfer": {
+            "gpu_ms": 0.9436812591552735,
+            "wall_ms": 0.9483181592077017
+          },
+          "marlin": {
+            "gpu_ms": 0.8909331512451172,
+            "wall_ms": 0.895569019485265
+          },
+          "order": [
+            "marlin",
+            "flashinfer"
+          ],
+          "round": 5
+        }
+      ],
+      "summary": {
+        "flashinfer": {
+          "gpu_ms": 0.9426467132568359,
+          "wall_ms": 0.9474158892408013
+        },
+        "marlin": {
+          "gpu_ms": 0.8909331512451172,
+          "wall_ms": 0.8956197183579206
+        },
+        "speedup_marlin_over_flashinfer": {
+          "gpu_ms": 0.945140038909117,
+          "wall_ms": 0.9453290033752898
+        }
+      }
+    }
+  }
+}

experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/phase4_aot_nojit_smoke_gpu6.json

@@ -0,0 +1,173 @@
+{
+  "compute_capability": [
+    12,
+    0
+  ],
+  "created_at": "2026-08-18T02:17:06.360422+00:00",
+  "cuda_visible_devices": null,
+  "device": "NVIDIA RTX 6000D",
+  "iterations": 1,
+  "rounds": 1,
+  "shapes": {
+    "decode": {
+      "case": {
+        "ep_rank": 7,
+        "ep_size": 32,
+        "hidden": 3584,
+        "inject_nonfinite": false,
+        "intermediate": 3072,
+        "local_experts": 28,
+        "name": "kimi_decode_m1_ep32",
+        "routing": "ep_mixed",
+        "tokens": 1,
+        "top_k": 16
+      },
+      "correctness": {
+        "flashinfer_vs_marlin": {
+          "cosine": 0.9996320605278015,
+          "finite": true,
+          "max_abs": 0.00978851318359375,
+          "mean_abs": 0.002149919280782342,
+          "nrmse": 0.02712450921535492,
+          "p95_abs": 0.00537109375,
+          "p99_abs": 0.007080078125,
+          "reference_max_abs": 0.37890625,
+          "reference_rms": 0.09998540580272675
+        },
+        "flashinfer_vs_reference": {
+          "cosine": 0.9996223449707031,
+          "finite": true,
+          "max_abs": 0.009765625,
+          "mean_abs": 0.00218405993655324,
+          "nrmse": 0.027483593672513962,
+          "p95_abs": 0.00537109375,
+          "p99_abs": 0.0072021484375,
+          "reference_max_abs": 0.37890625,
+          "reference_rms": 0.10006484389305115
+        },
+        "marlin_vs_reference": {
+          "cosine": 0.9999805688858032,
+          "finite": true,
+          "max_abs": 0.00390625,
+          "mean_abs": 0.00045798509381711483,
+          "nrmse": 0.0062673031352460384,
+          "p95_abs": 0.0009765625,
+          "p99_abs": 0.001953125,
+          "reference_max_abs": 0.37890625,
+          "reference_rms": 0.10006484389305115
+        }
+      },
+      "rounds": [
+        {
+          "flashinfer": {
+            "gpu_ms": 0.5889919996261597,
+            "wall_ms": 0.6109171081334352
+          },
+          "marlin": {
+            "gpu_ms": 0.5362880229949951,
+            "wall_ms": 0.6202820222824812
+          },
+          "order": [
+            "marlin",
+            "flashinfer"
+          ],
+          "round": 1
+        }
+      ],
+      "summary": {
+        "flashinfer": {
+          "gpu_ms": 0.5889919996261597,
+          "wall_ms": 0.6109171081334352
+        },
+        "marlin": {
+          "gpu_ms": 0.5362880229949951,
+          "wall_ms": 0.6202820222824812
+        },
+        "speedup_marlin_over_flashinfer": {
+          "gpu_ms": 0.9105183488661707,
+          "wall_ms": 1.015329271392741
+        }
+      }
+    },
+    "prefill": {
+      "case": {
+        "ep_rank": 7,
+        "ep_size": 32,
+        "hidden": 3584,
+        "inject_nonfinite": false,
+        "intermediate": 3072,
+        "local_experts": 28,
+        "name": "kimi_prefill_m128_ep32",
+        "routing": "ep_mixed",
+        "tokens": 128,
+        "top_k": 16
+      },
+      "correctness": {
+        "flashinfer_vs_marlin": {
+          "cosine": 0.9996392726898193,
+          "finite": true,
+          "max_abs": 0.0126953125,
+          "mean_abs": 0.0021022639703005552,
+          "nrmse": 0.02686966024339199,
+          "p95_abs": 0.005126953125,
+          "p99_abs": 0.0068359375,
+          "reference_max_abs": 0.462890625,
+          "reference_rms": 0.09850681573152542
+        },
+        "flashinfer_vs_reference": {
+          "cosine": 0.9996151924133301,
+          "finite": true,
+          "max_abs": 0.013671875,
+          "mean_abs": 0.00217337254434824,
+          "nrmse": 0.02774309180676937,
+          "p95_abs": 0.00537109375,
+          "p99_abs": 0.00689697265625,
+          "reference_max_abs": 0.46484375,
+          "reference_rms": 0.09856347739696503
+        },
+        "marlin_vs_reference": {
+          "cosine": 0.9999709129333496,
+          "finite": true,
+          "max_abs": 0.00390625,
+          "mean_abs": 0.000566127069760114,
+          "nrmse": 0.007645560894161463,
+          "p95_abs": 0.00146484375,
+          "p99_abs": 0.001953125,
+          "reference_max_abs": 0.46484375,
+          "reference_rms": 0.09856347739696503
+        }
+      },
+      "rounds": [
+        {
+          "flashinfer": {
+            "gpu_ms": 0.9108160138130188,
+            "wall_ms": 0.9314471390098333
+          },
+          "marlin": {
+            "gpu_ms": 0.813152015209198,
+            "wall_ms": 0.8787710685282946
+          },
+          "order": [
+            "marlin",
+            "flashinfer"
+          ],
+          "round": 1
+        }
+      ],
+      "summary": {
+        "flashinfer": {
+          "gpu_ms": 0.9108160138130188,
+          "wall_ms": 0.9314471390098333
+        },
+        "marlin": {
+          "gpu_ms": 0.813152015209198,
+          "wall_ms": 0.8787710685282946
+        },
+        "speedup_marlin_over_flashinfer": {
+          "gpu_ms": 0.8927730769741712,
+          "wall_ms": 0.943447063955196
+        }
+      }
+    }
+  }
+}

experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/phase4_mscan_ep32_5rounds_gpu6.json

@@ -0,0 +1,573 @@
+{
+  "compute_capability": [
+    12,
+    0
+  ],
+  "created_at": "2026-08-18T02:53:56.080340+00:00",
+  "cuda_visible_devices": "0",
+  "device": "NVIDIA RTX 6000D",
+  "iterations": 30,
+  "rounds": 5,
+  "shapes": {
+    "prefill": {
+      "case": {
+        "ep_rank": 7,
+        "ep_size": 32,
+        "hidden": 3584,
+        "inject_nonfinite": false,
+        "intermediate": 3072,
+        "local_experts": 28,
+        "name": "kimi_prefill_m128_ep32",
+        "routing": "ep_mixed",
+        "tokens": 128,
+        "top_k": 16
+      },
+      "correctness": {
+        "flashinfer_vs_marlin": {
+          "cosine": 0.9996392726898193,
+          "finite": true,
+          "max_abs": 0.0126953125,
+          "mean_abs": 0.0021022639703005552,
+          "nrmse": 0.02686966024339199,
+          "p95_abs": 0.005126953125,
+          "p99_abs": 0.0068359375,
+          "reference_max_abs": 0.462890625,
+          "reference_rms": 0.09850681573152542
+        },
+        "flashinfer_vs_reference": {
+          "cosine": 0.9996151924133301,
+          "finite": true,
+          "max_abs": 0.013671875,
+          "mean_abs": 0.00217337254434824,
+          "nrmse": 0.02774309180676937,
+          "p95_abs": 0.00537109375,
+          "p99_abs": 0.00689697265625,
+          "reference_max_abs": 0.46484375,
+          "reference_rms": 0.09856347739696503
+        },
+        "marlin_vs_reference": {
+          "cosine": 0.9999709129333496,
+          "finite": true,
+          "max_abs": 0.00390625,
+          "mean_abs": 0.000566127069760114,
+          "nrmse": 0.007645560894161463,
+          "p95_abs": 0.00146484375,
+          "p99_abs": 0.001953125,
+          "reference_max_abs": 0.46484375,
+          "reference_rms": 0.09856347739696503
+        }
+      },
+      "rounds": [
+        {
+          "flashinfer": {
+            "gpu_ms": 0.9099744160970052,
+            "wall_ms": 0.9106991346925497
+          },
+          "marlin": {
+            "gpu_ms": 0.8695210774739583,
+            "wall_ms": 0.871785699079434
+          },
+          "order": [
+            "marlin",
+            "flashinfer"
+          ],
+          "round": 1
+        },
+        {
+          "flashinfer": {
+            "gpu_ms": 0.910638936360677,
+            "wall_ms": 0.91134087027361
+          },
+          "marlin": {
+            "gpu_ms": 0.8706239700317383,
+            "wall_ms": 0.8713232663770517
+          },
+          "order": [
+            "flashinfer",
+            "marlin"
+          ],
+          "round": 2
+        },
+        {
+          "flashinfer": {
+            "gpu_ms": 0.928394635518392,
+            "wall_ms": 0.9290497982874513
+          },
+          "marlin": {
+            "gpu_ms": 0.8668832143147787,
+            "wall_ms": 0.8675735288610061
+          },
+          "order": [
+            "marlin",
+            "flashinfer"
+          ],
+          "round": 3
+        },
+        {
+          "flashinfer": {
+            "gpu_ms": 0.9538229624430339,
+            "wall_ms": 0.9544791265701255
+          },
+          "marlin": {
+            "gpu_ms": 0.8697290420532227,
+            "wall_ms": 0.8704239269718528
+          },
+          "order": [
+            "flashinfer",
+            "marlin"
+          ],
+          "round": 4
+        },
+        {
+          "flashinfer": {
+            "gpu_ms": 0.9259029388427734,
+            "wall_ms": 0.9265734348446131
+          },
+          "marlin": {
+            "gpu_ms": 0.8654357274373372,
+            "wall_ms": 0.8660859350735942
+          },
+          "order": [
+            "marlin",
+            "flashinfer"
+          ],
+          "round": 5
+        }
+      ],
+      "summary": {
+        "flashinfer": {
+          "gpu_ms": 0.9259029388427734,
+          "wall_ms": 0.9265734348446131
+        },
+        "marlin": {
+          "gpu_ms": 0.8695210774739583,
+          "wall_ms": 0.8704239269718528
+        },
+        "speedup_marlin_over_flashinfer": {
+          "gpu_ms": 0.9391060779661385,
+          "wall_ms": 0.9394009090254388
+        }
+      }
+    },
+    "prefill_m2048": {
+      "case": {
+        "ep_rank": 7,
+        "ep_size": 32,
+        "hidden": 3584,
+        "inject_nonfinite": false,
+        "intermediate": 3072,
+        "local_experts": 28,
+        "name": "kimi_prefill_m2048_ep32",
+        "routing": "ep_mixed",
+        "tokens": 2048,
+        "top_k": 16
+      },
+      "correctness": {
+        "flashinfer_vs_marlin": {
+          "cosine": 0.9996377229690552,
+          "finite": true,
+          "max_abs": 0.013671875,
+          "mean_abs": 0.0021010376513004303,
+          "nrmse": 0.026925193145871162,
+          "p95_abs": 0.005126953125,
+          "p99_abs": 0.0068359375,
+          "reference_max_abs": 0.5546875,
+          "reference_rms": 0.09821589291095734
+        },
+        "flashinfer_vs_reference": {
+          "cosine": 0.9996229410171509,
+          "finite": true,
+          "max_abs": 0.013671875,
+          "mean_abs": 0.002144214231520891,
+          "nrmse": 0.027460230514407158,
+          "p95_abs": 0.00537109375,
+          "p99_abs": 0.0068359375,
+          "reference_max_abs": 0.55859375,
+          "reference_rms": 0.09827093780040741
+        },
+        "marlin_vs_reference": {
+          "cosine": 0.9999799132347107,
+          "finite": true,
+          "max_abs": 0.005859375,
+          "mean_abs": 0.00045405104174278677,
+          "nrmse": 0.006360496394336224,
+          "p95_abs": 0.0009765625,
+          "p99_abs": 0.001953125,
+          "reference_max_abs": 0.55859375,
+          "reference_rms": 0.09827093780040741
+        }
+      },
+      "rounds": [
+        {
+          "flashinfer": {
+            "gpu_ms": 4.343468729654948,
+            "wall_ms": 4.344448462749521
+          },
+          "marlin": {
+            "gpu_ms": 9.073473103841145,
+            "wall_ms": 9.075486830746135
+          },
+          "order": [
+            "marlin",
+            "flashinfer"
+          ],
+          "round": 1
+        },
+        {
+          "flashinfer": {
+            "gpu_ms": 4.388905843098958,
+            "wall_ms": 4.389951967944701
+          },
+          "marlin": {
+            "gpu_ms": 9.065022786458334,
+            "wall_ms": 9.065993401842812
+          },
+          "order": [
+            "flashinfer",
+            "marlin"
+          ],
+          "round": 2
+        },
+        {
+          "flashinfer": {
+            "gpu_ms": 4.355183919270833,
+            "wall_ms": 4.35632496761779
+          },
+          "marlin": {
+            "gpu_ms": 9.069696044921875,
+            "wall_ms": 9.070580204327902
+          },
+          "order": [
+            "marlin",
+            "flashinfer"
+          ],
+          "round": 3
+        },
+        {
+          "flashinfer": {
+            "gpu_ms": 4.35906016031901,
+            "wall_ms": 4.360084631480277
+          },
+          "marlin": {
+            "gpu_ms": 9.070074462890625,
+            "wall_ms": 9.071069831649462
+          },
+          "order": [
+            "flashinfer",
+            "marlin"
+          ],
+          "round": 4
+        },
+        {
+          "flashinfer": {
+            "gpu_ms": 4.34208984375,
+            "wall_ms": 4.343056639966865
+          },
+          "marlin": {
+            "gpu_ms": 9.072443644205729,
+            "wall_ms": 9.073465705538789
+          },
+          "order": [
+            "marlin",
+            "flashinfer"
+          ],
+          "round": 5
+        }
+      ],
+      "summary": {
+        "flashinfer": {
+          "gpu_ms": 4.355183919270833,
+          "wall_ms": 4.35632496761779
+        },
+        "marlin": {
+          "gpu_ms": 9.070074462890625,
+          "wall_ms": 9.071069831649462
+        },
+        "speedup_marlin_over_flashinfer": {
+          "gpu_ms": 2.0825927517681464,
+          "wall_ms": 2.082275748269046
+        }
+      }
+    },
+    "prefill_m512": {
+      "case": {
+        "ep_rank": 7,
+        "ep_size": 32,
+        "hidden": 3584,
+        "inject_nonfinite": false,
+        "intermediate": 3072,
+        "local_experts": 28,
+        "name": "kimi_prefill_m512_ep32",
+        "routing": "ep_mixed",
+        "tokens": 512,
+        "top_k": 16
+      },
+      "correctness": {
+        "flashinfer_vs_marlin": {
+          "cosine": 0.9996380805969238,
+          "finite": true,
+          "max_abs": 0.01318359375,
+          "mean_abs": 0.0021038446575403214,
+          "nrmse": 0.026911206543445587,
+          "p95_abs": 0.005126953125,
+          "p99_abs": 0.0068359375,
+          "reference_max_abs": 0.5,
+          "reference_rms": 0.0984053835272789
+        },
+        "flashinfer_vs_reference": {
+          "cosine": 0.9996137619018555,
+          "finite": true,
+          "max_abs": 0.012939453125,
+          "mean_abs": 0.002174302702769637,
+          "nrmse": 0.027790850028395653,
+          "p95_abs": 0.00537109375,
+          "p99_abs": 0.0069580078125,
+          "reference_max_abs": 0.498046875,
+          "reference_rms": 0.09845462441444397
+        },
+        "marlin_vs_reference": {
+          "cosine": 0.9999691247940063,
+          "finite": true,
+          "max_abs": 0.005859375,
+          "mean_abs": 0.0005840405938215554,
+          "nrmse": 0.007868881337344646,
+          "p95_abs": 0.00146484375,
+          "p99_abs": 0.001953125,
+          "reference_max_abs": 0.498046875,
+          "reference_rms": 0.09845462441444397
+        }
+      },
+      "rounds": [
+        {
+          "flashinfer": {
+            "gpu_ms": 1.81865603129069,
+            "wall_ms": 1.8193970046316583
+          },
+          "marlin": {
+            "gpu_ms": 2.5967445373535156,
+            "wall_ms": 2.5989100647469363
+          },
+          "order": [
+            "marlin",
+            "flashinfer"
+          ],
+          "round": 1
+        },
+        {
+          "flashinfer": {
+            "gpu_ms": 1.7714890797932943,
+            "wall_ms": 1.7721137730404735
+          },
+          "marlin": {
+            "gpu_ms": 2.600511932373047,
+            "wall_ms": 2.6012318286423883
+          },
+          "order": [
+            "flashinfer",
+            "marlin"
+          ],
+          "round": 2
+        },
+        {
+          "flashinfer": {
+            "gpu_ms": 1.7895263671875,
+            "wall_ms": 1.7902158355961244
+          },
+          "marlin": {
+            "gpu_ms": 2.6133761088053387,
+            "wall_ms": 2.6140614645555615
+          },
+          "order": [
+            "marlin",
+            "flashinfer"
+          ],
+          "round": 3
+        },
+        {
+          "flashinfer": {
+            "gpu_ms": 1.7714154561360678,
+            "wall_ms": 1.7720839319129784
+          },
+          "marlin": {
+            "gpu_ms": 2.6104395548502604,
+            "wall_ms": 2.611126432505747
+          },
+          "order": [
+            "flashinfer",
+            "marlin"
+          ],
+          "round": 4
+        },
+        {
+          "flashinfer": {
+            "gpu_ms": 1.7820064544677734,
+            "wall_ms": 1.7826298717409372
+          },
+          "marlin": {
+            "gpu_ms": 2.6109214782714845,
+            "wall_ms": 2.6116128312423825
+          },
+          "order": [
+            "marlin",
+            "flashinfer"
+          ],
+          "round": 5
+        }
+      ],
+      "summary": {
+        "flashinfer": {
+          "gpu_ms": 1.7820064544677734,
+          "wall_ms": 1.7826298717409372
+        },
+        "marlin": {
+          "gpu_ms": 2.6104395548502604,
+          "wall_ms": 2.611126432505747
+        },
+        "speedup_marlin_over_flashinfer": {
+          "gpu_ms": 1.4648878225471482,
+          "wall_ms": 1.4647608423366596
+        }
+      }
+    },
+    "prefill_m8192": {
+      "case": {
+        "ep_rank": 7,
+        "ep_size": 32,
+        "hidden": 3584,
+        "inject_nonfinite": false,
+        "intermediate": 3072,
+        "local_experts": 28,
+        "name": "kimi_prefill_m8192_ep32",
+        "routing": "ep_mixed",
+        "tokens": 8192,
+        "top_k": 16
+      },
+      "correctness": {
+        "flashinfer_vs_marlin": {
+          "cosine": 0.9996379613876343,
+          "finite": true,
+          "max_abs": 0.0146484375,
+          "mean_abs": 0.0021007475443184376,
+          "nrmse": 0.0269180778414011,
+          "p95_abs": 0.005126953125,
+          "p99_abs": 0.0068359375,
+          "reference_max_abs": 0.5703125,
+          "reference_rms": 0.09823954105377197
+        },
+        "flashinfer_vs_reference": {
+          "cosine": 0.9996234178543091,
+          "finite": true,
+          "max_abs": 0.0146484375,
+          "mean_abs": 0.0021431385539472103,
+          "nrmse": 0.027442343533039093,
+          "p95_abs": 0.00537109375,
+          "p99_abs": 0.0068359375,
+          "reference_max_abs": 0.5703125,
+          "reference_rms": 0.09829506278038025
+        },
+        "marlin_vs_reference": {
+          "cosine": 0.9999799132347107,
+          "finite": true,
+          "max_abs": 0.005859375,
+          "mean_abs": 0.0004543001123238355,
+          "nrmse": 0.006363501772284508,
+          "p95_abs": 0.0009765625,
+          "p99_abs": 0.001953125,
+          "reference_max_abs": 0.5703125,
+          "reference_rms": 0.09829506278038025
+        }
+      },
+      "rounds": [
+        {
+          "flashinfer": {
+            "gpu_ms": 16.167036946614584,
+            "wall_ms": 16.168367435845237
+          },
+          "marlin": {
+            "gpu_ms": 34.38465983072917,
+            "wall_ms": 34.38796966802329
+          },
+          "order": [
+            "marlin",
+            "flashinfer"
+          ],
+          "round": 1
+        },
+        {
+          "flashinfer": {
+            "gpu_ms": 16.178935750325522,
+            "wall_ms": 16.179970768280327
+          },
+          "marlin": {
+            "gpu_ms": 34.393994140625,
+            "wall_ms": 34.39523999889692
+          },
+          "order": [
+            "flashinfer",
+            "marlin"
+          ],
+          "round": 2
+        },
+        {
+          "flashinfer": {
+            "gpu_ms": 16.147456868489584,
+            "wall_ms": 16.148571266482275
+          },
+          "marlin": {
+            "gpu_ms": 34.362894694010414,
+            "wall_ms": 34.36435852975895
+          },
+          "order": [
+            "marlin",
+            "flashinfer"
+          ],
+          "round": 3
+        },
+        {
+          "flashinfer": {
+            "gpu_ms": 16.169014485677085,
+            "wall_ms": 16.169893369078636
+          },
+          "marlin": {
+            "gpu_ms": 34.378727213541666,
+            "wall_ms": 34.37981543441614
+          },
+          "order": [
+            "flashinfer",
+            "marlin"
+          ],
+          "round": 4
+        },
+        {
+          "flashinfer": {
+            "gpu_ms": 16.138995361328124,
+            "wall_ms": 16.140034631825984
+          },
+          "marlin": {
+            "gpu_ms": 34.34004313151042,
+            "wall_ms": 34.341279161162674
+          },
+          "order": [
+            "marlin",
+            "flashinfer"
+          ],
+          "round": 5
+        }
+      ],
+      "summary": {
+        "flashinfer": {
+          "gpu_ms": 16.167036946614584,
+          "wall_ms": 16.168367435845237
+        },
+        "marlin": {
+          "gpu_ms": 34.378727213541666,
+          "wall_ms": 34.37981543441614
+        },
+        "speedup_marlin_over_flashinfer": {
+          "gpu_ms": 2.126470504586844,
+          "wall_ms": 2.126362823632778
+        }
+      }
+    }
+  }
+}

experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/phase4_nsys_flashinfer_decode_iteration_5_cuda_api_sum_nvtx=flashinfer_decode_iteration_5.csv

@@ -0,0 +1,4 @@
+Time (%),Total Time (ns),Num Calls,Avg (ns),Med (ns),Min (ns),Max (ns),StdDev (ns),Name
+87.0,139383,10,13938.3,11355.5,9370,21685,4468.6,cudaLaunchKernelExC
+7.2,11500,12,958.3,916.0,880,1324,122.6,cudaGetDriverEntryPointByVersion
+5.8,9349,12,779.1,705.5,565,1516,256.7,cuTensorMapEncodeTiled

experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/phase4_nsys_flashinfer_prefill_iteration_5_cuda_api_sum_nvtx=flashinfer_prefill_iteration_5.csv

@@ -0,0 +1,4 @@
+Time (%),Total Time (ns),Num Calls,Avg (ns),Med (ns),Min (ns),Max (ns),StdDev (ns),Name
+85.2,127306,10,12730.6,10292.0,8707,19421,4222.6,cudaLaunchKernelExC
+7.8,11727,12,977.3,641.0,518,3906,959.1,cuTensorMapEncodeTiled
+7.0,10386,12,865.5,818.5,790,1158,114.1,cudaGetDriverEntryPointByVersion

experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/phase4_nsys_marlin_decode_iteration_5_cuda_api_sum_nvtx=marlin_decode_iteration_5.csv

@@ -0,0 +1,5 @@
+Time (%),Total Time (ns),Num Calls,Avg (ns),Med (ns),Min (ns),Max (ns),StdDev (ns),Name
+61.2,71109,5,14221.8,14981.0,12385,15556,1523.4,cudaLaunchKernel
+24.7,28632,2,14316.0,14316.0,13516,15116,1131.4,cuLaunchKernelEx
+11.3,13082,1,13082.0,13082.0,13082,13082,0.0,cudaLaunchKernelExC
+2.8,3307,5,661.4,627.0,593,747,67.7,cuKernelGetName

experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/phase4_nsys_marlin_decode_iteration_5_cuda_gpu_kern_sum_nvtx=marlin_decode_iteration_5.csv

@@ -0,0 +1,7 @@
+Time (%),Total Time (ns),Instances,Avg (ns),Med (ns),Min (ns),Max (ns),StdDev (ns),Name
+94.4,109374,2,54687.0,54687.0,38783,70591,22491.7,"void sglang::device::marlin_moe::Marlin<__nv_bfloat16, (long)562949953487106, (long)2814749767106568, (int)128, (int)1, (int)8, (int)4, (bool)1, (int)4, (int)2, (bool)0, (bool)1, (bool)0>(const int4 *, const int4 *, int4 *, int4 *, const int4 *, const int4 *, const unsigned short *, const int4 *, const int *, const int *, const int *, const int *, const float *, int, bool, bool, int, int, int, int, int *, bool, bool, bool, int)"
+1.6,1888,1,1888.0,1888.0,1888,1888,0.0,"void sglang::topk_sum_kernel<(int)256, (bool)1>(sglang::TopkSumParams)"
+1.2,1344,1,1344.0,1344.0,1344,1344,0.0,_moe_align_small_numel_kernel
+1.1,1312,1,1312.0,1312.0,1312,1312,0.0,_situ_and_mul_kernel
+1.1,1248,2,624.0,624.0,608,640,22.6,"void at::native::vectorized_elementwise_kernel<(int)4, at::native::FillFunctor<c10::BFloat16>, std::array<char *, (unsigned long)1>>(int, T2, T3)"
+0.6,704,1,704.0,704.0,704,704,0.0,"void at::native::vectorized_elementwise_kernel<(int)4, at::native::FillFunctor<int>, std::array<char *, (unsigned long)1>>(int, T2, T3)"

experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/phase4_nsys_marlin_prefill_iteration_5_cuda_api_sum_nvtx=marlin_prefill_iteration_5.csv

@@ -0,0 +1,5 @@
+Time (%),Total Time (ns),Num Calls,Avg (ns),Med (ns),Min (ns),Max (ns),StdDev (ns),Name
+75.4,96744,7,13820.6,14257.0,10984,15791,1900.7,cudaLaunchKernel
+11.6,14830,1,14830.0,14830.0,14830,14830,0.0,cuLaunchKernelEx
+9.7,12416,1,12416.0,12416.0,12416,12416,0.0,cudaLaunchKernelExC
+3.3,4295,7,613.6,582.0,475,798,114.7,cuKernelGetName

experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/phase4_nsys_marlin_prefill_iteration_5_cuda_gpu_kern_sum_nvtx=marlin_prefill_iteration_5.csv

@@ -0,0 +1,8 @@
+Time (%),Total Time (ns),Instances,Avg (ns),Med (ns),Min (ns),Max (ns),StdDev (ns),Name
+94.1,709879,3,236626.3,181534.0,180702,347643,96144.2,"void sglang::device::marlin_moe::Marlin<__nv_bfloat16, (long)562949953487106, (long)2814749767106568, (int)256, (int)4, (int)16, (int)4, (bool)0, (int)4, (int)2, (bool)0, (bool)1, (bool)0>(const int4 *, const int4 *, int4 *, int4 *, const int4 *, const int4 *, const unsigned short *, const int4 *, const int *, const int *, const int *, const int *, const float *, int, bool, bool, int, int, int, int, int *, bool, bool, bool, int)"
+2.2,16671,1,16671.0,16671.0,16671,16671,0.0,_situ_and_mul_kernel
+1.7,13120,2,6560.0,6560.0,4896,8224,2353.3,"void at::native::vectorized_elementwise_kernel<(int)4, at::native::FillFunctor<c10::BFloat16>, std::array<char *, (unsigned long)1>>(int, T2, T3)"
+1.1,8448,1,8448.0,8448.0,8448,8448,0.0,"void sglang::topk_sum_kernel<(int)256, (bool)1>(sglang::TopkSumParams)"
+0.5,4064,1,4064.0,4064.0,4064,4064,0.0,"void moe_align_block_size_kernel<int>(const T1 *, int *, int *, int *, int, int, unsigned long, int *, bool, bool, int, int)"
+0.2,1696,1,1696.0,1696.0,1696,1696,0.0,"void count_and_sort_expert_tokens_kernel<int>(const T1 *, int *, int *, unsigned long, bool)"
+0.1,768,1,768.0,768.0,768,768,0.0,"void at::native::vectorized_elementwise_kernel<(int)4, at::native::FillFunctor<int>, std::array<char *, (unsigned long)1>>(int, T2, T3)"

experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/phase4_nsys_stats_cuda_api_sum.csv

@@ -0,0 +1,9 @@
+Time (%),Total Time (ns),Num Calls,Avg (ns),Med (ns),Min (ns),Max (ns),StdDev (ns),Name
+45.0,1657471,110,15067.9,13500.5,8707,37211,5252.2,cudaLaunchKernelExC
+27.8,1024642,60,17077.4,15240.0,10648,99929,11361.2,cudaLaunchKernel
+11.0,404869,1,404869.0,404869.0,404869,404869,0.0,cudaDeviceSynchronize
+7.7,283828,15,18921.9,16005.0,12865,41329,8082.9,cuLaunchKernelEx
+3.3,121690,120,1014.1,915.0,790,4516,414.6,cudaGetDriverEntryPointByVersion
+2.9,105105,120,875.9,712.0,518,3906,507.8,cuTensorMapEncodeTiled
+1.2,43805,60,730.1,650.0,468,2774,328.4,cuKernelGetName
+1.1,40880,1,40880.0,40880.0,40880,40880,0.0,cuProfilerStart

experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/phase4_nsys_stats_nvtx_gpu_proj_sum.csv

@@ -0,0 +1,21 @@
+Range,Style,Total Proj Time (ns),Total Range Time (ns),Range Instances,Proj Avg (ns),Proj Med (ns),Proj Min (ns),Proj Max (ns),Proj StdDev (ns),Total GPU Ops,Avg GPU Ops,Avg Range Lvl,Avg Num Child
+:marlin_decode_iteration_1,PushPop,1646632,2030238,1,1646632.0,1646632.0,1646632,1646632,0.0,8,8.0,0.0,0.0
+:flashinfer_prefill_iteration_1,PushPop,1499497,1236905,1,1499497.0,1499497.0,1499497,1499497,0.0,10,10.0,0.0,0.0
+:marlin_prefill_iteration_1,PushPop,1400491,1514107,1,1400491.0,1400491.0,1400491,1400491,0.0,9,9.0,0.0,0.0
+:marlin_decode_iteration_2,PushPop,1317581,1499161,1,1317581.0,1317581.0,1317581,1317581,0.0,8,8.0,0.0,0.0
+:flashinfer_decode_iteration_1,PushPop,1304108,1433573,1,1304108.0,1304108.0,1304108,1304108,0.0,10,10.0,0.0,0.0
+:marlin_prefill_iteration_2,PushPop,1252429,1302983,1,1252429.0,1252429.0,1252429,1252429,0.0,9,9.0,0.0,0.0
+:marlin_prefill_iteration_3,PushPop,1181582,1220545,1,1181582.0,1181582.0,1181582,1181582,0.0,9,9.0,0.0,0.0
+:marlin_decode_iteration_3,PushPop,1171887,1334668,1,1171887.0,1171887.0,1171887,1171887,0.0,8,8.0,0.0,0.0
+:marlin_prefill_iteration_4,PushPop,1161231,1184045,1,1161231.0,1161231.0,1161231,1161231,0.0,9,9.0,0.0,0.0
+:marlin_prefill_iteration_5,PushPop,1156238,1170223,1,1156238.0,1156238.0,1156238,1156238,0.0,9,9.0,0.0,0.0
+:marlin_decode_iteration_5,PushPop,1147087,1295194,1,1147087.0,1147087.0,1147087,1147087,0.0,8,8.0,0.0,0.0
+:flashinfer_prefill_iteration_2,PushPop,1132207,1082480,1,1132207.0,1132207.0,1132207,1132207,0.0,10,10.0,0.0,0.0
+:marlin_decode_iteration_4,PushPop,1126735,1274663,1,1126735.0,1126735.0,1126735,1126735,0.0,8,8.0,0.0,0.0
+:flashinfer_decode_iteration_2,PushPop,1081231,1167266,1,1081231.0,1081231.0,1081231,1081231,0.0,10,10.0,0.0,0.0
+:flashinfer_prefill_iteration_3,PushPop,1067216,1023953,1,1067216.0,1067216.0,1067216,1067216,0.0,10,10.0,0.0,0.0
+:flashinfer_prefill_iteration_4,PushPop,1003985,967242,1,1003985.0,1003985.0,1003985,1003985,0.0,10,10.0,0.0,0.0
+:flashinfer_decode_iteration_3,PushPop,991185,1054475,1,991185.0,991185.0,991185,991185,0.0,10,10.0,0.0,0.0
+:flashinfer_prefill_iteration_5,PushPop,955122,904533,1,955122.0,955122.0,955122,955122,0.0,10,10.0,0.0,0.0
+:flashinfer_decode_iteration_4,PushPop,924402,977971,1,924402.0,924402.0,924402,924402,0.0,10,10.0,0.0,0.0
+:flashinfer_decode_iteration_5,PushPop,900691,947759,1,900691.0,900691.0,900691,900691,0.0,10,10.0,0.0,0.0

experiments/pro6000/kimi3_pro6000_sglang_sm120_flashinfer_mxfp4/results/phase4_nsys_stats_nvtx_sum.csv

@@ -0,0 +1,21 @@
+Time (%),Total Time (ns),Instances,Avg (ns),Med (ns),Min (ns),Max (ns),StdDev (ns),Style,Range
+8.2,2030238,1,2030238.0,2030238.0,2030238,2030238,0.0,PushPop,:marlin_decode_iteration_1
+6.1,1514107,1,1514107.0,1514107.0,1514107,1514107,0.0,PushPop,:marlin_prefill_iteration_1
+6.1,1499161,1,1499161.0,1499161.0,1499161,1499161,0.0,PushPop,:marlin_decode_iteration_2
+5.8,1433573,1,1433573.0,1433573.0,1433573,1433573,0.0,PushPop,:flashinfer_decode_iteration_1
+5.4,1334668,1,1334668.0,1334668.0,1334668,1334668,0.0,PushPop,:marlin_decode_iteration_3
+5.3,1302983,1,1302983.0,1302983.0,1302983,1302983,0.0,PushPop,:marlin_prefill_iteration_2
+5.3,1295194,1,1295194.0,1295194.0,1295194,1295194,0.0,PushPop,:marlin_decode_iteration_5
+5.2,1274663,1,1274663.0,1274663.0,1274663,1274663,0.0,PushPop,:marlin_decode_iteration_4
+5.0,1236905,1,1236905.0,1236905.0,1236905,1236905,0.0,PushPop,:flashinfer_prefill_iteration_1
+5.0,1220545,1,1220545.0,1220545.0,1220545,1220545,0.0,PushPop,:marlin_prefill_iteration_3
+4.8,1184045,1,1184045.0,1184045.0,1184045,1184045,0.0,PushPop,:marlin_prefill_iteration_4
+4.8,1170223,1,1170223.0,1170223.0,1170223,1170223,0.0,PushPop,:marlin_prefill_iteration_5
+4.7,1167266,1,1167266.0,1167266.0,1167266,1167266,0.0,PushPop,:flashinfer_decode_iteration_2
+4.4,1082480,1,1082480.0,1082480.0,1082480,1082480,0.0,PushPop,:flashinfer_prefill_iteration_2
+4.3,1054475,1,1054475.0,1054475.0,1054475,1054475,0.0,PushPop,:flashinfer_decode_iteration_3
+4.2,1023953,1,1023953.0,1023953.0,1023953,1023953,0.0,PushPop,:flashinfer_prefill_iteration_3
+4.0,977971,1,977971.0,977971.0,977971,977971,0.0,PushPop,:flashinfer_decode_iteration_4
+3.9,967242,1,967242.0,967242.0,967242,967242,0.0,PushPop,:flashinfer_prefill_iteration_4
+3.8,947759,1,947759.0,947759.0,947759,947759,0.0,PushPop,:flashinfer_decode_iteration_5
+3.7,904533,1,904533.0,904533.0,904533,904533,0.0,PushPop,:flashinfer_prefill_iteration_5