Kernel Benchmarking Tales

benchmark 는 의외로 어렵다. 코드를 짜고 시간만 재면 끝날 것 같지만 — 실제로는 같은 코드가 셋팅에 따라 30%, 50%, 200% 다르게 나온다. 그래서 “이 새 커널이 빠르다” 의 주장이 거의 모두 — 측정의 함정에 빠진 채로 발표된다.

강의의 두 출발 질문.

1. 왜 같은 코드가 다르게 측정되는가 — warmup, cache, clock, jitter, 컴파일러 — 어느 것이 결정적 변수인가?
2. 어떻게 정직한 측정을 짤 것인가 — methodology 의 표준이 무엇이고, NVBench 같은 도구가 자동으로 잡아주는 것이 무엇인가?

본 노트는 transcript 실패로 — Georgii 의 공개 talk (CppCon, GTC), NVBench 의 README 와 benchmarking philosophy 문서, 그리고 보다 일반적인 GPU benchmarking 의 community 지식 (Sebastian Aaltonen 의 blog, Andreas Köpf 의 PMPP 강의) 으로 재구성.

cold call 은 — context init, kernel JIT, cache 첫 fill 이 섞여 매우 느림. warm call 만이 production 시간을 반영한다. 그래서 처음 N 회는 측정 전 “워밍업”으로 버린다.

# 표준 워밍업 패턴 (PyTorch)
import torch

# 1. warmup
for _ in range(10):
    out = kernel(in_data)
torch.cuda.synchronize()

# 2. measure
start = torch.cuda.Event(enable_timing=True)
end   = torch.cuda.Event(enable_timing=True)

start.record()
for _ in range(1000):
    out = kernel(in_data)
end.record()
torch.cuda.synchronize()

per_call_ms = start.elapsed_time(end) / 1000

benchmark 의 가장 미묘한 함정. 같은 입력 데이터를 반복 측정하면 — 두번째 호출부터 데이터가 L2 cache 에 있음. HBM 을 안 읽음. memory-bound 워크로드가 compute-bound 처럼 측정. 실전에서는 매 호출이 fresh data 라 — measurement 와 production 의 throughput 이 다름.

# NVBench 의 cache flush 패턴 (의사코드)
state.exec(
    [&](nvbench::launch& launch) {
        // flush L2 by reading large dummy buffer
        flush_l2(launch.get_stream());
        kernel<<<..., launch.get_stream()>>>(...);
    });
// NVBench 가 자동으로 매 call 마다 flush

또 다른 미묘한 자리 — SM (Streaming Multiprocessor) 의 internal state. 같은 kernel 을 반복 호출하면 — instruction cache 도 hit. 이건 production 에서도 hit 하는 게 정상이라 보존. 그러나 다른 kernel 들이 섞여 도는 시나리오에서는 instruction cache miss 가 일어남. micro-benchmark 가 over-optimistic.

한 번 측정한 시간 — 99% 정확하지 않다. clock jitter, OS scheduler, thermal, neighbour process — 모두 작은 noise 를 더함. 그래서 N 회 측정의 분포 로 봐야 한다. 평균만 보지 말고 — min, p50 (median), p99, std.

두 코드를 비교할 때 — “A 가 10.2ms, B 가 9.8ms 라서 B 가 4% 빠르다” 같은 결론은 위험. 만약 std 가 ±0.4ms 라면 두 분포가 겹친다 → 차이 유의미하지 않을 수 있음. 적어도 신뢰구간 겹침을 확인해야 함.

micro-benchmark — 한 kernel 의 측정. cache, launch overhead, isolated env. 깊은 분석에 좋음. macro-benchmark — 전체 학습/추론의 wall-clock. 진짜 production 시간. 두 측정의 결과가 일치하지 않을 수 있음.

benchmark harness 의 일을 손으로 짜는 것은 — 매번 같은 함정을 만든다. NVBench 같은 framework 가 그 함정을 자동 처리. Georgii 의 큰 contribution.

새 kernel 을 paper / blog / repo 에 발표할 때 — measurement 가 충분히 정직한가의 체크리스트. 강의의 가장 실용적 결론.

1. warmup 5+ 회 — 첫 cold call 은 측정에 안 들어감.
2. 측정 100+ 회 — 통계 처리. median 또는 trimmed mean.
3. std / 신뢰구간 보고 — 평균만 발표하지 말 것.
4. baseline 동일 환경 측정 — 자기 코드만 잘 측정하고 비교 baseline 은 default 셋팅이면 안 됨. 같은 GPU, 같은 driver, 같은 측정 도구.
5. cache flush 또는 fresh input — memory-bound kernel 은 cache 잔재가 결정적.
6. GPU clock 명시 — boost 인지 base 인지. base lock 권장.
7. compile flags 명시 — nvcc 의 optimization level. 다르게 컴파일하면 다른 결과.
8. dead-code 검증 — kernel 출력을 어디든 사용. 컴파일러가 통째로 제거 안 했는지.
9. 여러 input shape — 한 shape 에서만 빠르고 다른 shape 에서는 안 그럴 수 있음. 표 또는 그래프로.
10. 다른 GPU 에서 재측정 — A100 에서 빠른 게 H100 에서 느릴 수 있음.
11. reproducibility — 코드 + 측정 스크립트 + 환경 (driver, CUDA 버전) 명시. 다른 사람이 재현 가능.

• Georgii 의 구체적 “tales” — 강의 제목이 “Tales” 인 만큼 구체적 일화가 여러 개 있을 가능성. 본 노트는 일반적 함정으로 재구성. 원본 영상 확인 필요.
• NVBench 의 정확한 자동화 디테일 — 어떤 알고리즘으로 N 을 결정하는지, cache flush 방법은 어떻게 구현되는지.
• “흥미로운 buggy 측정” 사례 — 강의에서 실제 “이 paper 의 측정이 깨졌다” 같은 구체적 사례를 들었을 가능성.
• HPC 분야의 베스트 프랙티스 — Georgii 가 NVIDIA 에서 cuDNN/CUB 를 다루므로 HPC kernel 의 표준 측정에 더 깊은 디테일 가능성.
• multi-GPU benchmark — NCCL 시간 포함, GPU 간 sync 의 타임라인. 강의에서 다뤘는지.
• cloud / shared GPU 의 문제 — neighbour process 의 영향 통제 방법. lambdalabs 같은 환경에서.
• Hopper 의 새 측정 도구 — H100 의 hardware counter 가 새로 노출된 자리들.
• energy 측정 — perf/watt — 측정의 또 다른 차원. 강의에서 언급되었는지.