== PP0 rank ==
  file: /data/hf_models/pp2tp4_profiles/A_prefill/pa-1788706112.6645906-TP-0-PP-0.trace.json.gz
  baseTimeNanoseconds present: True
  span 35.176s  GPU-busy 34.065s (96.8%)  kernels 48052
  cudaMemcpy/Set GPU time: 598.4ms   cudaLaunchKernel calls: 4606 (421.1ms CPU)
  NCCL total: 33021.3ms (96.9% of busy, 93.9% of span)
    SendRecv       22601.8ms  n=170
    AllReduce      10419.5ms  n=2686
  top kernels (of GPU-busy):
     66.3%    22601.8ms  n=170     ncclDevKernel_SendRecv(ncclDevKernelArgsStorage
     30.6%    10419.5ms  n=2686    ncclDevKernel_AllReduce_Sum_bf16_RING_LL(ncclDevKernelArgsStorage
     24.9%     8483.5ms  n=1326    void sparse_mla_prefill_kernel
     16.6%     5661.4ms  n=11322   void cutlass::Kernel2
     11.9%     4044.2ms  n=2448    _ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalINS1_17GroupProblemShap
      2.6%      877.3ms  n=1224    void tensorrt_llm::kernels::cutlass_kernels::finalizeMoeRoutingKernel
      2.4%      819.2ms  n=408     void deep_gemm::sm120_fp8_mqa_logits
      2.0%      677.7ms  n=2618    kernel_cutlass_kernel_flashinfernormkernelsfused_add_rmsnormFusedAddRMSNormKerne
      1.8%      612.9ms  n=1224    void tensorrt_llm::kernels::cutlass_kernels::expandInputRowsKernel
      1.8%      598.4ms  n=1800    ::
      0.8%      271.0ms  n=1326    void sglang::concat_mla_absorb_q_kernel
      0.8%      261.2ms  n=2047    void at::native::vectorized_elementwise_kernel
      0.7%      236.9ms  n=1224    void tensorrt_llm::kernels::cutlass_kernels::doActivationKernel
      0.6%      221.1ms  n=408     (anonymous namespace)::topk_transform_prefill_kernel((anonymous namespace)::Fast
      0.2%       54.1ms  n=2652    kernel_cutlass_kernel_flashinfernormkernelsrmsnormRMSNormKernel_object_at__tenso
  forwards: 5 (gap>4.0ms split)
    wall ms: mean 6831.0  p50 1977.3  min 21.5  max 17604.3
    small forwards (<30% of p50): 1  (warmup/mixed steps)
    GPU-covered fraction inside main forwards: mean 99.4%
    inter-forward gap ms: mean 255.2  p50 243.9  p90 355.2  max 355.2

== PP1 rank ==
  file: /data/hf_models/pp2tp4_profiles/A_prefill/pa-1788706112.6645906-TP-0-PP-1.trace.json.gz
  baseTimeNanoseconds present: True
  span 35.181s  GPU-busy 34.609s (98.4%)  kernels 49218
  cudaMemcpy/Set GPU time: 600.3ms   cudaLaunchKernel calls: 5240 (1209.2ms CPU)
  NCCL total: 13796.9ms (39.9% of busy, 39.2% of span)
    AllReduce       9968.0ms  n=2652
    SendRecv        3590.7ms  n=170
    AG_RS            238.2ms  n=136
  top kernels (of GPU-busy):
     28.8%     9968.0ms  n=2652    ncclDevKernel_AllReduce_Sum_bf16_RING_LL(ncclDevKernelArgsStorage
     24.6%     8505.8ms  n=1326    void sparse_mla_prefill_kernel
     15.4%     5323.7ms  n=11235   void cutlass::Kernel2
     10.4%     3590.7ms  n=170     ncclDevKernel_SendRecv(ncclDevKernelArgsStorage
      8.5%     2930.6ms  n=2652    _ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalINS1_17GroupProblemShap
      2.7%      949.5ms  n=1326    void tensorrt_llm::kernels::cutlass_kernels::finalizeMoeRoutingKernel
      2.0%      697.6ms  n=2686    kernel_cutlass_kernel_flashinfernormkernelsfused_add_rmsnormFusedAddRMSNormKerne
      1.9%      668.9ms  n=1326    void tensorrt_llm::kernels::cutlass_kernels::expandInputRowsKernel
      1.7%      600.3ms  n=1687    ::
      1.0%      355.1ms  n=324     void deep_gemm::sm120_fp8_mqa_logits
      0.8%      278.8ms  n=2169    void at::native::vectorized_elementwise_kernel
      0.8%      270.7ms  n=1326    void sglang::concat_mla_absorb_q_kernel
      0.7%      245.0ms  n=1326    void tensorrt_llm::kernels::cutlass_kernels::doActivationKernel
      0.7%      238.2ms  n=136     ncclDevKernel_AllGather_RING_LL(ncclDevKernelArgsStorage
      0.5%      171.8ms  n=324     (anonymous namespace)::topk_transform_prefill_kernel((anonymous namespace)::Fast
  forwards: 10 (gap>4.0ms split)
    wall ms: mean 3482.4  p50 3786.0  min 4.3  max 9922.5
    small forwards (<30% of p50): 4  (warmup/mixed steps)
    GPU-covered fraction inside main forwards: mean 99.5%
    inter-forward gap ms: mean 39.7  p50 7.5  p90 277.8  max 277.8

== cross-stage bins (1.0ms, total 35.18s) ==
  both busy :    33.77s   96.0%
  only PP0  :     0.38s    1.1%   <- PP1 bubble
  only PP1  :     1.03s    2.9%   <- PP0 bubble
  neither   :     0.01s    0.0%   <- sched/CPU idle
  pipeline efficiency: both/(any busy) = 96.0%
