analyze_trace.py 4.2 KB

1234567891011121314151617181920212223242526272829303132333435363738394041424344454647484950515253545556575859606162636465666768697071
  1. import gzip, json, sys, collections, re
  2. path=sys.argv[1]
  3. opener = gzip.open if path.endswith('.gz') else open
  4. with opener(path,'rt') as f: data=json.load(f)
  5. ev=[e for e in data['traceEvents'] if e.get('ph')=='X']
  6. print("events:",len(ev))
  7. cats=collections.Counter(e.get('cat') for e in ev); print("cats:",dict(cats))
  8. gpu=[e for e in ev if e.get('cat') in ('kernel','gpu_memcpy','gpu_memset','gpu_user_annotation')]
  9. kern=[e for e in ev if e.get('cat') in ('kernel','gpu_memcpy','gpu_memset')]
  10. kern.sort(key=lambda e:e['ts'])
  11. cpu=[e for e in ev if e.get('cat') in ('cpu_op','cuda_runtime','user_annotation','python_function','cuda_driver')]
  12. t0=kern[0]['ts']; t1=max(e['ts']+e['dur'] for e in kern)
  13. print(f"GPU span: {(t1-t0)/1000:.1f} ms")
  14. # step boundaries: use cudaGraphLaunch runtime calls or user annotations
  15. launches=[e for e in ev if e.get('cat')=='cuda_runtime' and 'GraphLaunch' in e['name']]
  16. print("cudaGraphLaunch count:",len(launches))
  17. ann=[e for e in ev if e.get('cat') in ('user_annotation','gpu_user_annotation')]
  18. annc=collections.Counter(e['name'] for e in ann)
  19. print("annotations:",annc.most_common(25))
  20. def cat_of(n):
  21. l=n.lower()
  22. if 'nccl' in l: return 'nccl'
  23. if 'memcpy' in l or 'memset' in l: return 'memcpy/set'
  24. if 'engram' in l or 'hash' in l: return 'engram'
  25. if any(k in l for k in ('moe','expert','grouped','group_gemm','fp4','mxfp4','trtllm','topk','router','routing','silu','swiglu')): return 'moe'
  26. if any(k in l for k in ('attn','attention','mla','flash','indexer','index','sparse','paged','decode_kernel','rope','rotary','kv_cache','set_kv')): return 'attention'
  27. if any(k in l for k in ('gemm','matmul','cutlass','fp8','blockwise','scaled_mm','wgmma','bmm','sm90','sm100','sm120','cublas','gemv')): return 'gemm'
  28. if any(k in l for k in ('norm','rms')): return 'norm'
  29. if any(k in l for k in ('elementwise','vectorized','copy_','fill','cat','index_select','gather','scatter','arange','reduce','softmax','sort','cumsum','where','argmax','sampl','triton_')): return 'elementwise/other'
  30. return 'other'
  31. # union busy time
  32. busy=0; cur_s=None; cur_e=None
  33. for e in kern:
  34. s,en=e['ts'],e['ts']+e['dur']
  35. if cur_e is None or s>cur_e:
  36. if cur_e is not None: busy+=cur_e-cur_s
  37. cur_s,cur_e=s,en
  38. else: cur_e=max(cur_e,en)
  39. busy+=cur_e-cur_s
  40. print(f"GPU busy (union): {busy/1000:.1f} ms = {100*busy/(t1-t0):.1f}% of span")
  41. bycat=collections.defaultdict(float); cnt=collections.Counter()
  42. byname=collections.defaultdict(float); cntn=collections.Counter()
  43. for e in kern:
  44. c=cat_of(e['name']); bycat[c]+=e['dur']; cnt[c]+=1; byname[e['name']]+=e['dur']; cntn[e['name']]+=1
  45. print("\n== GPU time by category (sum of kernel durations, ms; count) ==")
  46. for c,v in sorted(bycat.items(),key=lambda kv:-kv[1]): print(f" {c:18s} {v/1000:9.1f} ms n={cnt[c]}")
  47. print("\n== top 40 kernels by total time ==")
  48. for n,v in sorted(byname.items(),key=lambda kv:-kv[1])[:40]: print(f" {v/1000:8.1f} ms n={cntn[n]:5d} avg={v/cntn[n]:8.1f}us {n[:110]}")
  49. # gaps
  50. gaps=[]; prev_e=kern[0]['ts']+kern[0]['dur']; prev_n=kern[0]['name']
  51. for e in kern[1:]:
  52. s=e['ts']
  53. if s-prev_e>300: gaps.append((s-prev_e,prev_e,prev_n,e['name']))
  54. prev_e=max(prev_e,s+e['dur']); prev_n=e['name'] if s+e['dur']>=prev_e else prev_n
  55. print(f"\n== GPU idle gaps >0.3ms: n={len(gaps)} total={sum(g[0] for g in gaps)/1000:.1f} ms ==")
  56. gb=collections.defaultdict(lambda:[0,0.0])
  57. for g in gaps:
  58. k=(g[2][:60],g[3][:60]); gb[k][0]+=1; gb[k][1]+=g[0]
  59. for k,v in sorted(gb.items(),key=lambda kv:-kv[1][1])[:15]: print(f" {v[1]/1000:8.1f} ms n={v[0]:4d} after [{k[0]}] before [{k[1]}]")
  60. # NCCL kernel stats
  61. nk=[e for e in kern if 'nccl' in e['name'].lower()]
  62. if nk:
  63. ds=sorted(e['dur'] for e in nk)
  64. print(f"\n== NCCL kernels: n={len(nk)} total={sum(ds)/1000:.1f} ms median={ds[len(ds)//2]:.0f}us p90={ds[int(len(ds)*.9)]:.0f}us max={ds[-1]:.0f}us")
  65. print(" names:",collections.Counter(e['name'][:80] for e in nk).most_common(6))
  66. # CPU-side top ops
  67. cpuop=collections.defaultdict(float); cpucnt=collections.Counter()
  68. for e in cpu:
  69. cpuop[(e['cat'],e['name'])]+=e['dur']; cpucnt[(e['cat'],e['name'])]+=1
  70. print("\n== top 25 CPU-side events by total time ==")
  71. for k,v in sorted(cpuop.items(),key=lambda kv:-kv[1])[:25]: print(f" {v/1000:8.1f} ms n={cpucnt[k]:5d} {k[0]}:{k[1][:90]}")