gpumode · 강의 아카이브
《GPU Mode》 L044 2025 · FEB · 22 Low priority · 보충 transcript · 실패 · 영상 직접 확인 필요

NVIDIA Profiling

NVIDIA 가 제공하는 profiling 도구 군 전체 — Nsight Systems · Nsight Compute · NVTX · Compute Sanitizer · Insight Graphics — 의 역할 분담과 진입 순서. transcript 가 빠졌으므로 본 노트는 도메인 지식 + L001 / L016 의 도구 어휘 위에서 정리한 보충 자료. 본문 주장은 영상 직접 확인 필요.

Nsight Systems Nsight Compute NVTX Compute Sanitizer CUPTI CLI vs GUI cloud GPU 한계 자동화 패턴
강의 번호
L044
스피커
미상 · 영상 메타데이터 누락
학습 우선순위
Low · L001 보충
상태
transcript fetch failed
§ 01강의가 풀려는 문제· profiler 지도가 왜 필요한가

NVIDIA 도구 6개 — 각자 다른 질문에 답한다

L001 이 “기본 프로파일러 4개” 의 진입 순서를 깔았다면, 이 강의는 NVIDIA 가 직접 제공하는 도구 군 전체의 지도를 본다. 같은 워크로드를 다른 도구로 보면 다른 정보가 잡힌다. 도구 선택이 곧 질문 선택이다.

강의 transcript 가 누락되어 정확한 발화 시퀀스는 영상 직접 확인이 필요하지만, 토픽이 가리키는 자리는 분명하다 — NVIDIA profiling 도구 전체의 역할 분담과 진입 순서를 한 시간 안에 깐다. PyTorch 사용자가 자기 워크로드의 어떤 질문에 어떤 도구를 들이대야 하는지의 의사결정 프레임.

강의의 인지적 frame · 추정

도구는 관찰 스케일 로 정렬된다. system 단위 (Nsys: CPU + GPU + 네트워크 동시) → kernel 단위 (NCU: 한 launch 안 stall) → line 단위 (NCU source view, PTX) → memory 단위 (Compute Sanitizer: race, OOB). 어느 스케일의 질문인지에 따라 도구가 결정된다.

“profiler 도구를 모르고 GPU 코드를 짠다는 건 어둠 속에서 못 박는 것과 같다 — 도구의 어휘가 곧 GPU 의 어휘다.”학습 노트 · L001 의 메시지 확장

본 노트의 작성 원칙 — transcript 가 없으므로 강의의 정확한 강조점은 보존하지 못한다. 대신 NVIDIA 공식 문서 + L001/L008/L016 의 강의 어휘 로 “이 강의가 거의 확실히 다뤘을 자리들” 을 정리한다. 모든 항목에 “영상 직접 확인 필요” 가 묵시적으로 붙어 있다.

§ 02NVIDIA 프로파일링 도구 군· Nsys · NCU · NVTX · Compute Sanitizer · Insight Graphics

도구 6개 — 각자 다른 질문에 답한다

현재 NVIDIA 가 제공하는 profiling 도구는 6개로 묶을 수 있다. 각 도구가 답하는 질문이 다르다는 사실 — 도구를 한 곳에 늘어놓고 보면 자명하다.

SYSTEM
Nsight Systems (nsys)
CPU + GPU + 네트워크 + memcpy + NCCL + Python 호출까지 한 timeline. 첫 진입점.
전체 학습 step 단위
KERNEL
Nsight Compute (ncu)
단일 커널 launch 안의 metric. memory hierarchy throughput, occupancy, warp stall, source-line 매핑까지.
단일 launch 단위
MARKER
NVTX
코드 안에서 의미 단위로 timeline 을 잘라준다. nvtx.range_push("attn_qk") 같이. Nsys / NCU 모두 인식.
사용자 정의 의미
CORRECT
Compute Sanitizer
CUDA 의 valgrind. race, OOB, uninitialized memory, sync error 검출. 성능보다 정확성.
디버깅용
GFX
Nsight Graphics
DirectX/Vulkan/OpenGL render pipeline. ML 워크로드 보다는 graphics 용. ML 사용자에겐 거의 안 쓰임.
graphics 전용
API
CUPTI
CUDA Profiling Tools Interface. 위 모든 도구가 이 위에 빌드. 직접 PyTorch profiler / Triton autotuner 가 사용.
SDK · 직접 호출 가능

이 6개의 관계 — CUPTI 가 가장 아래의 SDK고, 그 위에 Nsys, NCU 가 두 큰 도구로 자리 잡는다. NVTX 는 두 도구 모두에 메타데이터를 주는 marker. Compute Sanitizer 는 별도 라인 (정확성). Nsight Graphics 는 ML 외 영역.

진입 순서 — 추정

새 워크로드 → 먼저 Nsys 로 system trace. CPU-GPU 사이 idle gap, memcpy 패턴, NCCL 호출 모양 본다. 여기서 한 커널이 dominant 라는 게 보이면 → NCU 로 그 커널을 zoom. metric 이 hint 를 주면 → fix → 다시 NCU 로 확인. 일주일 뒤 → 새 코드 → 다시 Nsys. 무한 루프가 아니라 spiral 이다.

§ 03NCU 의 metric taxonomy· memory · pipeline · launch · stall

한 화면에 800+ metric — 어디부터 봐야 하는가

NCU 가 한 커널에 대해 보여주는 metric 수는 800 개를 넘는다. 모두 다 볼 수는 없다. NVIDIA 가 제공하는 section file 이 카테고리별로 묶어준다.

SpeedOfLightpeak 대비 현재 — compute throughput vs memory throughput. 가장 먼저 본다. 어느 쪽이 bottleneck 인지 1초 안에 결론.
MemoryWorkloadAnalysisL1 / L2 / device memory 각 계층의 traffic. coalescing, sector hit rate, bank conflict. L043 이 매 단계 본 화면.
SchedulerStatisticswarp 가 발행한 instruction 수 / 발행 가능한 cycle 수. 발행률이 낮으면 stall. 어떤 stall 인지의 분포가 다음 section.
WarpStateStatisticsstall reason 분포. Long Scoreboard / LG Throttle / Short Scoreboard / Math Pipe / Barrier / Selected. L043 § 08 의 어휘 그대로.
InstructionStatsinstruction 종류별 count. integer math, FP math, MMA, load/store, sync. fix 후 instruction count 가 감소하는지 검증.
LaunchStatsgrid · block 모양, register/thread, SMEM/block, occupancy. tail effect 의 진단 자료.
SourceCounterssource line / SASS instruction 별 stall sample. “이 한 줄이 70% 의 stall을 만든다” 를 보여준다.
Occupancyregister / SMEM / launch shape 가 occupancy 에 미치는 영향. theoretical vs achieved 차이의 원인.

이 8개 section 의 진입 순서가 곧 NCU 사용 워크플로 — SpeedOfLight 로 큰 그림Memory or Pipeline 둘 중 어느 쪽이 bottleneckWarpStateStatistics 로 stall 카테고리 확인Source view 로 책임 라인 식별.

NCU CLI vs GUI

CLI 로 빠르게 metric 확인 (ncu --section SpeedOfLight ./app) → 흥미로운 자리가 있으면 trace export (--set full -o trace.ncu-rep) → GUI 에서 같은 trace 를 열어 source view 까지. CI 에서는 CLI 만 사용 가능 — JSON export 로 metric regression 검사 가능.

§ 04Nsight Systems · system trace· CPU · GPU · NCCL · stream

“한 학습 step 의 한 장 사진” — 모든 일이 한 timeline 에 같이 나온다

Nsys 의 강점은 scope. NCU 가 한 커널 launch 안을 깊이 보는 도구라면, Nsys 는 여러 process / GPU / CPU thread 를 한 timeline 에 같이 그린다. distributed training 의 진단 도구.

10⁰ s
step 단위
전체 학습 step 의 wall-clock
10⁻³ s
Nsys 의 자리
ms 단위 사건 — 커널 / memcpy / NCCL / Python
10⁻⁶ s
NCU 의 자리
µs 단위 — 한 커널 launch 안
10⁻⁹ s
SASS 단위
ns 단위 — instruction 사이

Nsys 가 잘 잡는 자리들.

  • CPU-GPU idle gap — Python 이 다음 work 를 발행하지 못해서 GPU 가 비는 자리. dataloader bottleneck 의 흔한 형태.
  • memcpy H↔D 가 dominant — 입력 transfer 가 커널보다 길면 커널 최적화 효과 없음.
  • NCCL collective 의 형태 — all-reduce 가 compute 와 overlap 되는지 안 되는지. ring vs hierarchical 의 패턴 차이.
  • multi-stream 활용 — stream 두 개를 띄웠는데 실제로는 serialize 되고 있는지.
  • kernel launch overhead — 작은 커널을 너무 많이 발행. CUDA Graph 의 fix 대상.
# 가장 작은 Nsys 사용 — CLI
nsys profile -t cuda,nvtx,cudnn,cublas \
  -o profile.qdrep \
  python train.py --steps=10

# 결과 .qdrep 또는 .nsys-rep 를 GUI 로 열기
# 또는 CLI 로 stats 추출:
nsys stats profile.nsys-rep \
  --report cuda_kernel_sum

# Distributed: 각 rank 별 trace 를 GUI 에서 같이 열기
# nsys 가 다중 process 를 자동으로 한 화면에 정렬한다
Nsys 와 NCU 가 같이 사는 법

Nsys → 어디가 dominant 한 커널인지 식별 → 그 커널만 NCU 로 다시 본다 (ncu --kernel-id ::regex:gemm:). 두 도구가 같은 NVTX marker 를 인식하므로 — with nvtx.range("attn"): 안만 측정하는 식의 zoom 이 가능.

§ 05NVTX 마커 활용· 코드 의미 단위로 trace 잘라보기

도구가 잡는 단위와 사람이 보는 단위를 일치시키는 작은 도구

기본 프로파일링은 커널 이름 단위로 trace 를 잘라준다. 하지만 사람이 보고 싶은 단위는 보통 의미 단위 — “attention forward”, “MLP block”, “optimizer step”. NVTX 가 이 매핑.

# PyTorch + NVTX
import torch.cuda.nvtx as nvtx

def forward(x):
    nvtx.range_push("attention")
    nvtx.range_push("qk_matmul")
    qk = q @ k.transpose(-1, -2)
    nvtx.range_pop()

    nvtx.range_push("softmax")
    p = qk.softmax(-1)
    nvtx.range_pop()

    nvtx.range_push("pv_matmul")
    out = p @ v
    nvtx.range_pop()
    nvtx.range_pop()  # attention

# 이제 Nsys / NCU 의 timeline 에 "attention" 박스가 직접 보인다
# Nsys CLI 로 NVTX 별 stats 도 추출 가능:
#   nsys stats --report nvtx_sum profile.nsys-rep
실전 패턴

학습 코드에 한 번 NVTX layer (attention / mlp / norm / optimizer) 를 깔면 — 모든 후속 profiling 이 그 의미 단위로 정렬된다. PyTorch 의 torch.profiler 도 자동으로 NVTX 를 emit (with_stack=True 와 함께). Triton 의 autotuner 도 NVTX 를 emit 해서 어떤 config 가 어디서 도는지 추적 가능.

§ 06자주 보는 hint 들· tail · scoreboard · uncoalesced

NCU 가 직접 박아주는 actionable hint 5+1 개

NCU 가 좋은 점은 — metric 만 dump 하는 게 아니라 “이걸 고치면 X% 빨라진다” 의 hint 까지 박아준다는 것. ML 사용자가 자주 만나는 hint 5개를 정리.

Tail effect
마지막 wave 가 underfilled. grid shape 이 SM 수의 배수가 아닐 때. fix: padding 으로 입력 크기를 SM 의 배수에 맞춤.
Long Scoreboard Stalls
memory load 의 latency 가 hide 안 됨. fix: SMEM tile, prefetching, register reuse.
Uncoalesced Global Loads
warp 안의 thread 가 메모리에서 멀리 떨어진 자리 읽음. fix: column-major 변환, layout 정리, vectorized load.
Low Achieved Occupancy
register 또는 SMEM 사용량이 너무 많아서 SM 위 동시 warp 수 부족. fix: tile 줄이기 or __launch_bounds__.
SMEM Bank Conflict
한 warp 의 thread 들이 같은 bank 의 다른 word 에 접근. fix: array layout padding (예: +1 stride).
Register Spilling
사용 register 가 한계 초과해서 local memory(=HBM) 로 spill. fix: tile 줄이기, fragment 수 줄이기, 코드 단순화.

L043 의 5단계 사다리가 이 hint 들의 직접적 응답이었다는 것을 다시 떠올리면 — uncoalesced → column-major fix, long scoreboard → SMEM tile, tail effect → register tile. 각 hint 가 곧 fix 의 trigger.

§ 07cloud GPU 에서의 한계· NCU 권한 · MIG · multi-tenant

왜 클라우드에서 NCU 가 종종 안 되는지

실용적인 함정. NCU 는 hardware performance counter 에 직접 접근한다. 보안 및 multi-tenant 제약 때문에 cloud 환경에서는 종종 막혀 있다.

multi-tenant counter 격리한 GPU 위 두 사용자가 있으면 counter 가 서로 leak 될 수 있음 → cloud provider 가 counter access 를 disable. NCU 명령은 돌지만 결과가 비어 나옴.
MIG (Multi-Instance GPU)A100 / H100 의 MIG 모드에서 일부 hardware counter 가 instance 별로 격리. Nsys 는 동작, NCU 의 일부 metric 은 제한.
권한 (sudo)NCU 의 일부 고급 metric 은 root 또는 GPU performance counter 권한이 필요 (nvidia-smi -q -d PERFORMANCE_COUNTERS). 환경 별 정책 다름.
대안자기 데스크탑 GPU (RTX) → NCU 에 거의 항상 full access. cloud 에서는 Nsys 가 더 현실적. PyTorch profiler + chrome trace 도 항상 동작.
실전

대형 모델 실험 자체는 cloud H100 에서 돌리되, kernel-level 최적화는 자기 데스크탑(예: RTX 4090) 에서 한다 — 같은 architecture (Ada Lovelace 는 Hopper 와 호환 가능한 SASS) 에서 NCU full access 로 hint 받고, fix 한 코드를 cloud 에 올린다. 이 두 환경의 분리가 일반적인 워크플로.

§ 08자동화 패턴· section file · CI 통합

profiling 이 한 번 만지고 끝이 아니라 — 매번 자동으로 도는 검사

실전 팀의 패턴 — profiling 결과를 사람이 매번 읽지 않는다. 핵심 metric 만 추출해서 CI 의 regression test 로.

# NCU 결과를 JSON 으로 export
ncu --set full --target-processes all \
    --export profile.ncu-rep \
    --csv \
    -k regex:my_kernel_name \
    python bench.py

ncu --import profile.ncu-rep --csv \
    --metrics gpu__time_duration.sum,sm__throughput.avg.pct_of_peak_sustained_elapsed \
    --csv > metrics.csv

# Python 후처리 — 임계치 검사
import pandas as pd
df = pd.read_csv("metrics.csv")
peak_pct = df["sm__throughput.avg.pct_of_peak_sustained_elapsed"].mean()
assert peak_pct > 60, f"regression: {peak_pct}%"
section file 직접 작성

NCU 의 기본 section 외에, 자기 워크로드에 특화된 section 을 작성 가능. ~/Documents/NVIDIA Nsight Compute/Sections/.section 파일. ML 팀에서 “attention 전용 section” 또는 “GEMM 전용 section” 을 깔아두는 패턴이 흔하다.

이 자동화의 가치 — 사람이 매번 NCU GUI 열고 화면 비교하는 대신, commit 별 metric 변화를 dashboard 에 자동 plot. regression 이 PR 단계에서 잡힌다.

§ 09docs map· 어디서 무엇을 찾는가

NVIDIA 공식 문서의 진입점 — 6개

실전에서 “이 hint 의 정확한 의미가 뭐지?” 또는 “이 metric 어떻게 측정되지?” 를 찾을 때 가는 자리들.

NCU User Guide
CLI 옵션, section / set 의 의미. 가장 자주 보는 자리.
NCU Kernel Profiling Guide
metric 의 의미, hint 가 동작하는 logic. roofline model 까지.
CUDA Best Practices Guide
coalescing, occupancy, tiling 의 정석. 본 강의의 hint 들의 출처.
PTX ISA
PTX instruction reference. NCU 의 source view 에서 PTX 를 보다가 의문이 들 때.
Nsys User Guide
CLI 옵션, NVTX 통합, multi-process trace.
CUPTI Reference
프레임워크 (PyTorch profiler, Triton autotuner) 가 직접 부르는 SDK. 자기 도구 만들 때.
§ 10기억할 메모와 코드· key takeaways

다시 열었을 때 5분 안에 손에 잡혀야 할 것

Nsys 먼저, NCU 다음
새 워크로드는 Nsys 로 system trace → dominant kernel 식별 → NCU 로 zoom.
NVTX 한 번 깔기
학습 코드에 attention / mlp / norm / optimizer 의 NVTX range 한 번 깔면 모든 후속 trace 가 의미 단위로 정렬.
SpeedOfLight 가 첫 화면
NCU 열면 가장 먼저 — peak 대비 어디까지 왔는지. compute vs memory 어느 쪽 bottleneck 인지.
cloud 에서 NCU 막힐 수 있음
대안: 자기 데스크탑 GPU 에서 kernel-level 최적화. cloud 에서는 Nsys + PyTorch profiler.
CI 자동화
ncu --csv 로 metric export → regression test. PR 단계에서 throughput 회귀 잡기.
Compute Sanitizer
정확성 검사 도구. race / OOB / uninit. 성능 도구가 아니라 디버거. 새 커널 짜고 한 번은 돌린다.
§ 11다른 강의로 이어지는 길· connections

이 도구 지도가 다음에 어디에 다시 등장하는지

§ 12열린 질문· open questions

transcript 누락으로 더 많은 자리가 비어 있다

강의 transcript 가 누락된 강의 — 본 노트의 모든 자리는 도메인 지식 + L001/L008/L016 의 어휘 기반의 추정. 영상 직접 확인 필요한 사항들.

검증 메모

본 노트 전체에 “영상 직접 확인 필요” 가 묵시적으로 붙어 있다. 모든 numerical 주장 (예: "metric 800개+", "SM 별 64KB") 은 NVIDIA 공식 문서 기준이며 강의에서 직접 발화된 수치는 아니다. 본 노트는 강의 노트라기보다 L044 의 토픽이 가리키는 자리에 대한 도메인 지식 보충으로 읽는 게 정확하다.

← Lecture 043 int8 tensorcore matmul for Turing Lecture 045 → Outperforming cuBLAS on H100 — pranjalssh