CUDA 성능 분석 도구 다른거 사용하자, (NVTX - 성능 분석용 태깅 도구) 실제 커널최적화의 경우 Ncu ( Nsight Compute)

Nsight Systems 는 타임라인 / API 레벨 중심

Nsight Compute 는 단일 커널 성능 분석 중심

 

 Nsight Systems ( Nsys)

타임라인 기반, 프로그램 전체 흐름, 스레드, 스트림, CUDA Graph, NVTX 범위 등을 보는 도구 현재 사용중

관찰 레벨

• 전체 실행 타임라인
• 커널/메모리 호출 순서
• 스트림 상의 동시성
• CPU,  GPU 상호 작용
• CUDA Graph 캡처 안정성
• NVTX 범위
• 병목 구간

 

Nsight Compute( Ncu )

커널 하나를 자세히 분석

GEMM 커널 최적화할 때 필수...!!

• SM occupancy
• warp stall reason
• 메모리 대역폭 사용률
• shared memory bank confilct
• register pressure
• scheduler issue slot
• Tensor Core 활용률 미친
• cp.async pipeline stall
• instruction-level timeline

 

 

CUPTI Raw API ( 직접 사용 ) - 할까 말까

Nsight 가 내부에서 사용하는 원천 API

• 커널 실행 시간
• 메모리 트래픽
• 그래프 런치 여부
• stream event 

실시간 직접 수집 가능, 추후 내부 profiler 를 직접 만들 때 CUPTI 기반으로 ...

 

실제 Nsight Compute 기반 gemm 내용 

4. Nsight Compute 기반 GEMM 커널 분석

ncu로 실행했을 때 일부 환경 문제(드라이버 불일치)와 경고가 있었지만,
핵심 metrics는 정상 수집됨.

4.1. 성능 지표 핵심

항목값의미

Duration	14.43 µs	tiny GEMM → 오버헤드 중심
Compute Throughput	0.87%	작은 workload라 compute 거의 사용 안됨
Memory Throughput	<1%	동일 이유
Achieved Occupancy	~17%	warp 수가 적어서 latency 숨기지 못함
Active Warps per SM	~8	Ampere 기준 매우 적은 수
Major stall	L1TEX 의존성 (scoreboard wait)	global load coalescing suboptimal

4.2. Nsight Compute 오류 원인

• Cuda driver is not compatible with Nsight Compute
  → 현재 ncu 버전과 NVIDIA 드라이버 버전 불일치
• ERR_NVGPUCTRPERM
  → Performance Counter 접근 권한 설정 필요
• Launch Statistics 항목 대부분 n/a
  → 드라이버/모듈 로딩 실패로 profiling pass 일부 무효

4.3. 분석 결론

• tiny GEMM(m=32,k=64,n=16) 기준에서 GPU가 거의 일을 하지 않으므로
  throughput, occupancy, IPC가 낮게 나오는 건 정상적인 구조적 현상.
• 다만 global load의 coalescing 문제는 실제 커널에서 확인해야 할 부분.

---

5. 환경 단계 정리

5.1. 현재 해결된 항목

• pyd/glue/layer 호출 체계
• NVTX instrumentation
• CUDA Graph capture/replay 안정성
• GEMM forward/backward correctness

5.2. 아직 해결해야 하는 항목

• Nsight Compute와 CUDA Driver 버전 호환성
• GPU Performance Counter 권한 설정
• 일부 Launch Statistics 수집 실패
• Python 종료 시 utf-8-sig 인코딩 에러 발생

---

6. 향후 개선 및 최적화 계획

아키텍처 레벨 → 커널 레벨 순서로 정리.

---

6.1. 아키텍처 / 런타임 레벨

(1) Nsight Compute 환경 정상화

• NVIDIA 드라이버 최신 버전으로 업데이트
• GPU Performance Counter 활성화
  (NVIDIA 제어판 → Developer Settings → Counters: 모든 사용자 허용)
• CUDA Toolkit에 포함된 ncu 버전과 맞춰 재실행

(2) 중간 규모 테스트 벤치 추가

• tiny GEMM → GPU 활용도 분석 불가
  → m=n=k=256, m=n=k=1024 수준 벤치 스크립트 추가
• 그래야 occupancy/stall/memory pattern이 의미 있게 나타남

(3) Nsight Systems + NVTX로 layer-level 프로파일링 자동화

• 각 layer forward/backward에 NVTX range 태깅
• Graph capture 구간에 대해 자동 validity check

---

6.2. 커널 레벨 최적화

(1) 메모리 coalescing 개선

• Source Counters가 global load uncoalescing을 감지
• suboptimal stride → thread/block indexing 재점검 필요
• warp당 연속 load(32 threads → 32 contiguous floats)를 보장하도록 구조화

(2) occupancy 확보

• warp 수 증가 → latency hiding 유리
• 점검 항목:
  • register pressure
  • shared memory usage
  • block size(BM/BN/BK tiling)
  • grid launch config(grid.x, block.x)

(3) cp.async 기반 double buffering (Ampere 이상)

• global→shared prefetch를 warp-level로 pipelining
• L1TEX scoreboard stall 44% → 크게 개선 가능

(4) MMA tensor core path 추가

• FP16/TF32 MMA warp tile (16×8×16 / 16×16×16)
• 커널 variant 추가: f32 path vs mma path

(5) epilogue fusion 정식화

• 현재 bias + activation fuse 커널을
  structured epilogue로 분리하여 여러 활성화/정규화 대응

---

6.3. 문서/테스트 고도화

(1) ops 단위 정식 smoke test

• forward/backward correctness (현재 OK)
• gradient check
• random input repeat/replay consistency

(2) Graph Capture Safety Check 리스트화

• 모든 ops에서 unsafe API(동기 memcpy, sync, malloc 등) 금지
• shim 매크로 기반 자동 검증

(3) NVTX 기반 layer-level 보고서 자동 생성

• 실행 후 “layer별 latency 리포트” 자동 프린트