제2부 CUDA 프로그래밍

2.14 프로파일링 II — Nsight Compute

2.14.1 이 장에서 배우는 것 / 선수 지식

2.13에서 nsys로 시스템 병목을 보고 핫 커널을 골랐다. 이 장은 그 커널 하나를 Nsight Compute(ncu) 로 해부해 "무엇에 바운드인가"를 판정한다 — SOL·메모리 차트·점유율·warp stall·루프라인. 그리고 측정 → 가설 → 수정 → 재측정 사이클을 한 커널에 끝까지 적용한다.

배우는 것:

• ncu로 커널 메트릭을 수집하고 Speed Of Light(SOL)·메모리 차트·점유율·루프라인으로 메모리 vs 연산 바운드를 판정한다.
• warp stall 사유로 latency 바운드를 진단한다.
• ncu replay 오버헤드를 절대시간으로 착각하지 않는다.

선수 지식: 2.13(nsys·시스템 병목·측정 규율), 1.5(워프 stall), 1.6(점유율), 1.10(루프라인·산술 강도). 1.10의 임계 강도 I* = 피크÷BW를 안다고 가정 — 이 장은 그 종이 계산을 도구로 확인한다.

환경 주의(S-5): 메트릭 이름·SOL 섹션은 GPU 세대(ncu는 Volta+)·도구 버전 의존. 출력은 예시(개략). 정확한 메트릭은 ncu --query-metrics·공식 문서(A.3).

2.14.2 수집 — ncu

ncu는 커널을 여러 번 재실행(replay) 하며 하드웨어 카운터를 수집한다 — 느리지만 점유율·SOL·메모리 처리량·루프라인을 다 준다.

ncu --set full -k fused_norm_kernel -c 1 -o ncu_report ./my_app   # 특정 커널 첫 1회, 풀 섹션
ncu --set basic -k gemm_tile_kernel -c 1 ./my_app                 # 핵심 섹션만(빠름)
ncu --metrics sm__throughput.avg.pct_of_peak_sustained_elapsed,\
gpu__dram_throughput.avg.pct_of_peak_sustained_elapsed -k softmax_kernel -c 1 ./my_app

• -k·-c로 반드시 범위를 좁힌다. 안 그러면 모든 커널을 모든 인스턴스에서 replay해 측정이 안 끝난다.
• --set full은 루프라인·SOL·메모리 차트까지(replay 많음), --set basic은 빠른 1차 진단.
• replay 동안 입력이 바뀌면 안 된다(같은 커널이 매번 같은 일). 누적·아토믹으로 상태가 변하는 커널은 해석 주의.
• GUI(ncu-ui)로 .ncu-rep를 열면 섹션별 차트, CLI는 ncu --import ... --page details.

2.14.3 Speed Of Light (SOL) — 바운드 판정의 핵심 한 줄

SOL 섹션은 이 커널이 두 천장(연산·메모리) 에 각각 몇 %나 닿았는지를 준다(1.10 루프라인의 실측판).

판정: Memory 88% ≫ Compute 18% → 메모리 바운드.

판정 규칙(개략):

• Memory % 높고(>60%) Compute % 낮음 → 메모리 바운드. 약은 코어가 아니라 융합·재사용·정밀도 축소로 산술 강도 I 높이기(1.10).
• Compute % 높고 Memory % 낮음 → 연산 바운드. 약은 텐서코어·정밀도·점유율(1.11·1.13·1.6).
• 둘 다 낮음(둘 다 <50%) → latency 바운드/비효율. 점유율 부족·warp stall·비병합·런치 오버헤드 의심(아래). 가장 흔하고 고칠 것이 많은 경우.

정직: SOL %는 "피크 대비"라 GPU 세대·정밀도·SOL 정의에 따라 다른 수가 나온다. 절대 % 자체보다 두 수의 상대 관계(어느 쪽이 천장에 가까운가)가 판정의 핵심.

2.14.4 메모리 차트·점유율·warp stall

2.14.4.1 메모리 차트 — 어느 단계가 막혔나

데이터가 SM ↔ L1/TEX ↔ L2 ↔ DRAM(HBM)을 어떻게 오갔는지, 각 링크의 처리량·적중률을 보여준다.

• sector/request 비율: 병합이 잘 되면 이상적(요청당 sector 적음), 비병합이면 sector가 부풀어 실효 대역폭↓(1.8). DRAM은 한가한데 L1/L2 트래픽만 폭증하면 접근 패턴 의심.
• L2 hit rate: 낮으면 재사용이 안 돼 HBM을 자꾸 다시 친다. 타일링(2.7)·레이아웃 점검.
• DRAM throughput %: HBM이 천장에 가까우면(SOL Memory와 일치) 실제 메모리 바운드 — 더 줄이려면 이동량 자체를 줄여야(융합·정밀도).

2.14.4.2 점유율 — 지연을 은닉할 워프가 충분한가

Theoretical Occupancy   50%   ← 자원(레지스터/공유메모리/블록크기)이 정한 상한
Achieved Occupancy      31%   ← 실제로 채워진 비율
Limiter:  Registers  (스레드당 64 reg → SM당 워프 수 제한)

• 이론 점유율이 낮다 → 자원이 제약. Limiter가 Registers면 -Xptxas -v로 레지스터·스필 확인(1.6), __launch_bounds__ 검토. 공유 메모리면 타일 크기, 블록 크기면 launch 구성.
• 이론은 높은데 달성이 낮다 → 실행 중 불균형(꼬리 효과·작업 분배 불균형) 또는 워프가 빨리 끝나 못 채움.
• 주의: 점유율은 수단이지 목표가 아니다(1.6). 점유율을 높이려 레지스터를 줄였다가 스필이 나면 오히려 느려진다. 메모리 바운드면 점유율 100%여도 천장은 HBM. SOL로 바운드부터 판정.

2.14.4.3 warp stall — 워프가 왜 멈춰 있나

둘 다 SOL이 낮은 latency 바운드 커널의 실제 원인이 여기서 나온다(1.5).

• Long Scoreboard 우세 → 글로벌 메모리 대기. 점유율↑·접근 병합·타일링·prefetch.
• Barrier 우세 → __syncthreads 대기(2.8). 동기화 줄이기·불균형 작업 재분배.
• MIO/LSU 우세 → 공유 메모리 뱅크 충돌(1.7)·로드 스토어 포화. 접근 패턴·벡터화.

2.14.4.4 루프라인 섹션

--set full은 루프라인 차트를 그린다. 1.10에서 종이로 찍던 점(x=I, y=달성 성능)을 실측으로 산출해 준다.

• 점이 경사선(BW×I) 위에 가까이 → 대역폭을 잘 쓰는 메모리 바운드(I를 높인다: 융합·재사용·정밀도).
• 점이 수평선 근처 → 연산 바운드(텐서코어·점유율).
• 점이 경사선보다도 한참 아래 → 대역폭조차 못 채우는 비효율(비병합·낮은 점유율·런치 오버헤드·분기). 1.10.5와 같은 진단을 도구로 확정.

2.14.5 워크스루 — 측정 → 가설 → 수정 → 재측정

입력 텐서에 정규화 + 활성화를 하는 element-wise 커널이 nsys에서 GPU 시간의 60%를 먹는 핫 커널로 잡혔다고 하자.

1. nsys로 핫 커널 확정: [Kernel Summary]에서 norm_act_kernel 60%. 타임라인에서 커널 사이 갭이 작고 SM이 빽빽 → 시스템 병목 아님, 커널 내부 문제 → ncu.
2. ncu로 SOL·루프라인: ncu --set full -k norm_act_kernel -c 1 ./app → SOL Memory 85%, Compute 12% → 메모리 바운드. 루프라인 점이 경사선 위 가까움 → 대역폭은 잘 쓰는 중. L2 hit 낮음 → 재사용 없음(element-wise라 당연).
3. 가설: SOL이 메모리 바운드라 말한다. 코어 튜닝은 무의미. 이 커널은 입력을 읽고 중간 텐서를 쓰고 다음 커널이 그것을 다시 읽는다 → 앞뒤 element-wise와 융합하면 중간 텐서 HBM 왕복이 사라진다(트래픽 ~1/3 → 시간 ~1/3). 또는 FP16/BF16(2B)으로 이동량 절반.
4. 수정 + 재측정: 융합 커널로 바꾼 뒤 같은 측정 반복. 커널 시간↓, Memory SOL은 여전히 높음(여전히 메모리 바운드)이지만 절대 이동 byte 감소 → 절대 시간 단축. nsys로 전체 반복 시간 단축 확인.

핵심: 재측정으로 가설을 확정하거나 기각. 시간이 기대만큼 안 줄었다면 — 융합이 레지스터를 늘려 점유율을 낮췄거나(occupancy 섹션), 비병합이 새로 생겼을 수 있다(memory 차트). 사이클을 다시 순환한다. ncu는 두 리포트를 baseline 비교해 수정 전후 메트릭 증감(Δ)을 색으로 보여준다.

2.14.6 바운드 판정 순서도 (체크포인트)

임의의 ncu 리포트를 받아 5분 안에 판정한다:

자가 점검: SOL Memory 10%, Compute 8%, 점유율 25%, Long Scoreboard 압도적이면? → 둘 다 낮음 = latency 바운드. 메모리 지연이 노출됐고 점유율이 낮아 지연을 은닉하지 못한다. "코어가 약하다/HBM이 약하다"가 아니라 워프가 부족하거나 접근이 비병합이라 판정.

2.14.7 트레이드오프 — 경량 대안

견해: 루틴은 이벤트 타이밍으로 "빨라졌나"를 낮은 비용으로 확인하고, 느려지면 nsys로 시스템부터, 핫 커널을 ncu --set basic → 필요하면 full로 좁힌다. 모든 커널에 --set full은 시간 낭비 — 80%를 먹는 한두 개만.

2.14.8 흔한 함정

• ncu replay 오버헤드를 절대시간으로 착각. ncu는 커널을 여러 번 replay하므로 ncu로 잰 "전체 실행 시간"은 의미가 없다(수십~수백 배 과대 측정됨). ncu는 커널 내부 메트릭·상대 비교용, 절대 시간·처리량은 nsys(또는 이벤트)로.
• 점유율 = 성능으로 착각. Achieved Occupancy 100%가 빠름을 보장 안 한다. 메모리 바운드면 천장은 HBM. SOL로 바운드부터 판정한 뒤 점유율을.
• -k·-c 미지정. 모든 커널을 모든 인스턴스에서 replay해 안 끝난다. 범위를 좁혀라.
• 한 입력 크기로 일반화. 같은 GEMM도 큰 배치는 연산 바운드, 디코딩(배치1 GEMV)은 메모리 바운드(1.10). 대표적 실제 크기로.

2.14.9 연습문제

연습 1 (SOL 판정). ncu 출력이 Memory SOL 82%, Compute SOL 15%, Achieved Occupancy 70%. 무엇에 바운드이며 1순위는? 점유율을 100%로 높이면 빨라지나?

연습 2 (둘 다 낮음). Memory SOL 12%, Compute SOL 9%, Occupancy 20%, Warp Stall은 Long Scoreboard 압도적. 진단명과 후보 원인·확인 항목 2가지.

연습 3 (ncu 시간 오해). 동료가 "ncu로 재니 이 커널이 200배 느려졌다, 수정이 망가졌다"고 한다. 무엇이 틀렸나? 절대 시간은 어떻게?

연습 4 (개방형). 2.7의 타일링 행렬곱을 nsys로 핫 여부 확인하고 ncu --set full로 SOL·점유율·루프라인을 수집해 바운드를 판정하라. 하나의 수정(타일 크기·블록 구성·정밀도)을 가설과 함께 적용하고 재측정해 가설이 맞았는지 보고하라.

2.14.10 요약 · 다음 장

• ncu(Nsight Compute) 는 커널 하나의 SOL·메모리 차트·점유율·warp stall·루프라인으로 왜 느린지를 답한다.
• 바운드 판정은 SOL 한 줄: Memory≫Compute = 메모리 바운드(I 높이기), Compute≫Memory = 연산 바운드(텐서코어·점유율), 둘 다 낮음 = latency 바운드(점유율·병합·stall로). 루프라인 점 위치로 교차 확인 — 1.10의 종이 계산을 도구로 확정.
• ncu는 메트릭·상대비교용, 절대시간은 nsys·이벤트로. 점유율 = 성능 아님 — SOL로 바운드부터.

제2부(CUDA)가 곧 끝난다. 다음 장: 2.15 CUDA 라이브러리 I — Thrust·CUB·CCCL. 직접 커널을 작성하기 전에 검증된 라이브러리를 먼저 고르는 길 — 표준 연산은 라이브러리가 거의 항상 이긴다.