Nsight Compute 프로파일링

어느 기둥이 깨졌는지 알려주는 CUDA 프로파일러

Nsight Compute는 NVIDIA의 per-kernel profiler. 도는 CUDA process에 attach해서 모든 커널 launch에 대해 상세 counter 캡처, 읽기 좋은 dashboard로 제시. BLAS 작업엔 가장 중요한 metric 여섯:

Metric	값 (RTX 4090에서 cuBLAS GEMM 4096³)	의미
Compute (SM) Throughput	80.27%	Tensor pipe 켜짐 — 목표 그거
Memory Throughput	42.54%	L2 도움; DRAM 절반 idle
DRAM Throughput	18.15%	GEMM이 compute-bound, BW-bound 아님 — 확인
L2 Hit Rate	93.64%	Tile size가 operand를 L2에 hot 유지
Achieved Occupancy	16.64%	SM당 ~8 warp 스케줄 (reg + smem 한계)
FMA pipeline busy	~73%	GEMM이 fused MUL-ADD instruction 지배

Internal 커널 duration: FP16-Tensor-Core path에서 2.63 ms ≈ 52 TF/s. Wall time ~2118 ms 인 이유는 Nsight가 instrumented profiling pass 40+ 주입. Nsight 안에서 perf 판단할 땐 항상 'internal' 시간 읽고, wall time 아냐.

그림 읽기

그 표가 'healthy' 모습. 커널이 compute-bound (80% SM), L2이 일 잘 함 (94% hit rate), FMA pipe가 limiting factor. 헤드룸 essentially 없음 — 라이브러리 기둥이 이미 max한 ~26%-from-perfect 커널.

Unhealthy 커널 모습:

SM Throughput < 30% + DRAM Throughput > 80% → bandwidth-bound. core 안 먹임.
SM Throughput < 30% + DRAM < 30% → idle / launch 오버헤드 지배. Block 너무 적거나 커널 너무 짧음.
L2 Hit Rate < 50% → tile size가 안 맞음. tile 줄이거나 주변 op fuse.
Occupancy < 10% → register spilling. --maxrregcount 시도 또는 더 적은 register 쓰게 재구성.

Code

CUDA binary에 Nsight Compute 돌리기·bash

# 위 표 metric 캡처:
ncu --set full ./gemm_bench --m 4096 --n 4096 --k 4096

# 더 가벼운 캡처 (BLAS-relevant metric만):
ncu --metrics sm__cycles_active.avg.pct_of_peak_sustained_elapsed,\
    dram__throughput.avg.pct_of_peak_sustained_elapsed,\
    lts__t_sector_hit_rate.pct,\
    sm__warps_active.avg.pct_of_peak_sustained_active \
    ./gemm_bench

# 'gemm_tiled' 이름 커널만 프로파일 (warmup launch 스킵):
ncu --kernel-name gemm_tiled ./gemm_bench

Programmatic Nsight start/stop range (launch 많을 때)·cuda

#include <nvtx3/nvToolsExt.h>

// 진짜 프로파일하고 싶은 region 표시
nvtxRangePush("hot_loop");
for (int i = 0; i < 100; ++i) {
    cublasSgemm(...);     // 진짜 워크로드
}
nvtxRangePop();

// 그 range만 캡처:
//   ncu --nvtx --nvtx-include 'hot_loop' ./your_app

Exercise

CUDA 박스에서 Track 6의 tiled GEMM을 ncu --set full로 프로파일. 위 표의 metric 여섯 캡처. 같은 shape에 cuBLAS 리포트랑 비교. 'SM Throughput'이랑 cuBLAS 80% 사이 격차가 정확히 최적화 여지. 그 다음 각 격차가 어느 기둥 반영하는지 봐: DRAM Throughput 80%인데 SM 30%면 더 나은 tiling 필요. SM 50% DRAM 20%면 Tensor Core path 필요.

어느 기둥이 깨졌는지 알려주는 CUDA 프로파일러

그림 읽기

Code

External links

Exercise

Progress

댓글 0