prof-b1-analysis.txt 3.9 KB

1234567891011121314151617181920212223242526272829303132333435363738394041424344454647
  1. decode steps analysed: 42
  2. step wall ms: mean=71.7 median=71.3 min=67.6 max=81.9
  3. GPU busy per step: mean=67.6 ms
  4. == per-step GPU time by category (ms): mean / median / max ; kernels per step ==
  5. bf16_gemm 29.3 / 29.1 / 35.5 n=420
  6. moe_grouped_gemm 21.0 / 20.8 / 23.6 n=86
  7. nccl 12.7 / 12.6 / 15.7 n=100
  8. hyperconn 2.2 / 2.1 / 3.2 n=344
  9. attention 2.1 / 2.1 / 2.8 n=346
  10. moe_aux 1.6 / 1.6 / 1.9 n=221
  11. other 1.3 / 1.3 / 1.6 n=612
  12. router_tiny_gemm 0.9 / 0.9 / 1.0 n=43
  13. quant 0.8 / 0.8 / 1.0 n=234
  14. norm 0.4 / 0.4 / 0.5 n=140
  15. memcpy 0.3 / 0.2 / 0.4 n=60
  16. fp8_triton_gemm 0.1 / 0.1 / 0.2 n=1
  17. engram 0.0 / 0.0 / 0.0 n=6
  18. SUM 72.6
  19. NCCL per step ms: mean=12.7 median=12.6 min=10.6 max=15.7
  20. NCCL kernel duration percentiles us: p10=27 p25=29 p50=37 p75=155 p90=379 p95=513 p99=749
  21. NCCL kernels <150us: 75% time share of >300us kernels: 55%
  22. prof-b1-rank0.trace.json.gz: engram-following NCCL: n=90 total=31.0 ms mean=345us median=284us max=1794us | other NCCL n=4395 total=555.3 ms | engram share 5%
  23. steps 44
  24. 23.76 ms/step n/step= 87.8 _ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalINS1_17GroupProblemShapeIN4cute5tupleIJlllEEEEENS1_10
  25. 16.98 ms/step n/step= 233.8 kernel_cutlass_kernel_flashinfergemmkernelsdense_blockscaled_gemm_sm120_b12xDenseGemmKernel_object_at__CopyAto
  26. 12.94 ms/step n/step= 91.9 ncclDevKernel_AllReduce_Sum_bf16_RING_LL(ncclDevKernelArgsStorage<4096ul>)
  27. 11.48 ms/step n/step= 46.0 void cutlass::Kernel2<cutlass_80_wmma_tensorop_bf16_s161616gemm_bf16_16x16_128x2_tn_align8>(cutlass_80_wmma_te
  28. 1.34 ms/step n/step= 87.8 _hc_mix_stats_partial_kernel
  29. 1.30 ms/step n/step= 40.9 void flashinfer::sparse_mla_sm120::sparse_mla_decode_dsv4_kernel<(ModelType)1, 64, 128, 64>(__nv_bfloat16 cons
  30. 0.87 ms/step n/step= 40.0 void sglang::tiny_n_gemm_kernel<sglang::GEMMTraitN<384u, 5120u, 8u, 32u>, 4u, float, true>(sglang::TinyGEMMPar
  31. 0.78 ms/step n/step= 233.8 kernel_cutlass_kernel_flashinferquantizationkernelsmxfp8_quantizeMXFP8QuantizeSwizzledKernel_object_at__tensor
  32. 0.59 ms/step n/step= 87.8 _hc_mix_reduce_sinkhorn_kernel
  33. 0.41 ms/step n/step= 43.9 _router_triton_kernel
  34. 0.39 ms/step n/step= 43.9 void tensorrt_llm::kernels::cutlass_kernels::expandInputRowsKernel<__nv_fp8_e4m3, __nv_fp8_e4m3, (tensorrt_llm
  35. 0.35 ms/step n/step= 4.0 ncclDevKernel_AllGather_RING_LL(ncclDevKernelArgsStorage<4096ul>)
  36. 0.34 ms/step n/step= 43.9 void tensorrt_llm::kernels::quantize_with_block_size<(tensorrt_llm::BlockScaleQuantizationType)2, __nv_bfloat1
  37. 0.34 ms/step n/step= 43.9 void sglang::silu_mul_clamp_kernel<__nv_bfloat16, true>(sglang::SiluAndMulClampParams)
  38. 0.27 ms/step n/step= 90.9 kernel_cutlass_kernel_flashinfernormkernelsrmsnormRMSNormKernel_object_at__tensorptrbf16gmemalign128oi64512051
  39. 0.27 ms/step n/step= 43.9 void tensorrt_llm::kernels::cutlass_kernels::finalizeMoeRoutingKernel<__nv_bfloat16, __nv_bfloat16, __nv_bfloa
  40. 0.25 ms/step n/step= 43.9 void tensorrt_llm::kernels::cutlass_kernels::doActivationKernel<__nv_fp8_e4m3, __nv_bfloat16, __nv_bfloat16, t
  41. 0.22 ms/step n/step= 3.0 void cutlass::Kernel2<cutlass_80_wmma_tensorop_bf16_s161616gemm_bf16_16x16_128x1_tn_align8>(cutlass_80_wmma_te
  42. 0.22 ms/step n/step= 97.0 void at::native::elementwise_kernel<128, 4, at::native::gpu_kernel_impl_nocast<at::native::direct_copy_kernel_
  43. 0.22 ms/step n/step= 40.0 void tensorrt_llm::kernels::cutlass_kernels::fusedBuildExpertMapsSortFirstTokenKernel<32, 6, 8>(int const*, in
  44. 0.21 ms/step n/step= 43.9 _page_split_kernel
  45. 0.21 ms/step n/step= 87.8 _mhc_post_split_h_kernel