== PP0 rank ==
  file: /data/hf_models/pp2tp4_profiles/D_128k/pd-1788708582.9504747-TP-0-PP-0.trace.json.gz
  baseTimeNanoseconds present: True
  span 19.900s  GPU-busy 19.311s (97.0%)  kernels 30279
  cudaMemcpy/Set GPU time: 212.3ms   cudaLaunchKernel calls: 8256 (1914.1ms CPU)
  NCCL total: 10898.6ms (56.4% of busy, 54.8% of span)
    SendRecv        6970.9ms  n=77
    AllReduce       3927.6ms  n=1264
    AG_RS              0.0ms  n=2
  top kernels (of GPU-busy):
     36.1%     6970.9ms  n=77      ncclDevKernel_SendRecv(ncclDevKernelArgsStorage
     28.0%     5397.8ms  n=576     void deep_gemm::sm120_fp8_mqa_logits
     20.3%     3927.6ms  n=1264    ncclDevKernel_AllReduce_Sum_bf16_RING_LL(ncclDevKernelArgsStorage
     17.3%     3347.8ms  n=507     void sparse_mla_prefill_kernel
     10.9%     2106.7ms  n=4986    void cutlass::Kernel2
      6.7%     1300.1ms  n=1152    _ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalINS1_17GroupProblemShap
      3.1%      608.0ms  n=576     (anonymous namespace)::topk_transform_prefill_kernel((anonymous namespace)::Fast
      2.5%      486.4ms  n=583     void at::native::vectorized_gather_kernel
      1.7%      334.3ms  n=576     void tensorrt_llm::kernels::cutlass_kernels::finalizeMoeRoutingKernel
      1.3%      257.0ms  n=1232    kernel_cutlass_kernel_flashinfernormkernelsfused_add_rmsnormFusedAddRMSNormKerne
      1.2%      232.6ms  n=576     void tensorrt_llm::kernels::cutlass_kernels::expandInputRowsKernel
      1.1%      212.3ms  n=1735    ::
      0.5%      102.2ms  n=624     void sglang::concat_mla_absorb_q_kernel
      0.5%       99.5ms  n=1635    void at::native::vectorized_elementwise_kernel
      0.5%       87.3ms  n=576     void tensorrt_llm::kernels::cutlass_kernels::doActivationKernel
  forwards: 7 (gap>4.0ms split)
    wall ms: mean 2769.3  p50 583.1  min 3.1  max 11507.6
    small forwards (<30% of p50): 3  (warmup/mixed steps)
    GPU-covered fraction inside main forwards: mean 99.1%
    inter-forward gap ms: mean 85.8  p50 77.5  p90 231.2  max 231.2

== PP1 rank ==
  file: /data/hf_models/pp2tp4_profiles/D_128k/pd-1788708582.9504747-TP-0-PP-1.trace.json.gz
  baseTimeNanoseconds present: True
  span 30.873s  GPU-busy 19.218s (62.2%)  kernels 36752
  cudaMemcpy/Set GPU time: 222.0ms   cudaLaunchKernel calls: 13111 (5216.2ms CPU)
  NCCL total: 6021.4ms (31.3% of busy, 19.5% of span)
    AllReduce       3757.5ms  n=1248
    SendRecv        2191.8ms  n=82
    AG_RS             72.1ms  n=69
  top kernels (of GPU-busy):
     23.1%     4429.8ms  n=1008    void deep_gemm::sm120_fp8_mqa_logits
     19.6%     3757.5ms  n=1248    ncclDevKernel_AllReduce_Sum_bf16_RING_LL(ncclDevKernelArgsStorage
     17.4%     3340.8ms  n=507     void sparse_mla_prefill_kernel
     11.4%     2191.8ms  n=82      ncclDevKernel_SendRecv(ncclDevKernelArgsStorage
     10.3%     1973.9ms  n=4939    void cutlass::Kernel2
      5.8%     1113.6ms  n=1248    _ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalINS1_17GroupProblemShap
      2.9%      547.9ms  n=1008    (anonymous namespace)::topk_transform_prefill_kernel((anonymous namespace)::Fast
      1.9%      361.8ms  n=624     void tensorrt_llm::kernels::cutlass_kernels::finalizeMoeRoutingKernel
      1.9%      355.7ms  n=1022    void at::native::vectorized_gather_kernel
      1.4%      264.8ms  n=1264    kernel_cutlass_kernel_flashinfernormkernelsfused_add_rmsnormFusedAddRMSNormKerne
      1.3%      252.9ms  n=624     void tensorrt_llm::kernels::cutlass_kernels::expandInputRowsKernel
      1.2%      222.0ms  n=2576    ::
      0.6%      106.7ms  n=2063    void at::native::vectorized_elementwise_kernel
      0.5%      102.4ms  n=624     void sglang::concat_mla_absorb_q_kernel
      0.5%       91.4ms  n=624     void tensorrt_llm::kernels::cutlass_kernels::doActivationKernel
  forwards: 4 (gap>4.0ms split)
    wall ms: mean 4847.5  p50 6553.0  min 5.1  max 12800.7
    small forwards (<30% of p50): 2  (warmup/mixed steps)
    GPU-covered fraction inside main forwards: mean 99.2%
    inter-forward gap ms: mean 3827.7  p50 553.0  p90 10924.9  max 10924.9

== cross-stage bins (1.0ms, total 30.87s) ==
  both busy :    18.79s   60.9%
  only PP0  :     0.59s    1.9%   <- PP1 bubble
  only PP1  :     0.56s    1.8%   <- PP0 bubble
  neither   :    10.93s   35.4%   <- sched/CPU idle
  pipeline efficiency: both/(any busy) = 94.2%
