Lecture 4 · Compute and Memory Basics

PMPP Ch.1–3 (L002) 까지는 “GPU 위에 일을 어떻게 분배하는가” 의 framing. 거기서 짠 커널이 빠른지 느린지는 측정해 볼 수는 있어도 왜 그런지 는 답할 수 없다. L004 가 그 자리를 메운다 — Ch.4(GPU 아키텍처)와 Ch.5(메모리 모델 + 성능)를 묶어 “이 커널은 compute 가 모자란가, memory bandwidth 가 모자란가, occupancy 가 모자란가” 의 첫 정량 진단을 가르친다.

강의가 답하는 질문은 셋.

1. SM 안에서 무엇이 무엇을 도는가 — FP32 unit, Tensor Core, warp scheduler, register file, shared memory.
2. 같은 데이터에 어떤 메모리 계층이 적용되는가 — 그리고 각 계층의 latency · bandwidth 가 대략 얼마인가.
3. roofline 위에서 내 커널이 어디 있는가 — peak FLOPs 의 70% 인지, peak bandwidth 의 70% 인지, 아니면 둘 다 못 채우고 있는지.

강의의 첫 정량 — 한 SM 안에 32개의 FP32 unit, 그 중 절반은 INT32 도 가능, 그리고 별도의 Tensor Core unit. 한 cycle 에 처리되는 thread 의 단위가 warp = 32 threads 이고, 한 SM 은 여러 warp 를 동시에 schedule 한다 (Ampere 기준 4 warp scheduler).

이 표의 두 가지 사실이 이 강의의 모든 디테일에 반복적으로 등장한다.

• warp = 32 threads, 한 cycle 에 한 instruction 을 lockstep 으로 실행 — 그래서 warp 안에서 분기 (if-else) 가 갈라지면 두 path 가 직렬로 돌면서 절반은 idle. 이게 warp divergence.
• register file 이 SM 당 256 KB — 한 thread 당 몇 register 를 쓰는지가 SM 위에 동시에 둘 수 있는 thread 수를 직접 제한. occupancy (§06) 의 출발점.

강의에서 Thomas 가 가장 못 박는 사실 — “CUDA 는 block 을 SM 들에 임의로 배정한다. 우리는 실행 순서를 모른다.” 이 한 문장이 GPU 의 확장성의 비밀이고, 동시에 학습자가 가장 자주 잊는 사실이다.

이 사실의 결과 두 가지.

• block 사이의 통신은 kernel boundary 가 유일 — block A 가 block B 의 결과를 보려면 두 개의 kernel launch 가 필요하다. “이 block 이 다른 block 의 결과를 기다린다” 같은 코드는 deadlock 의 길.
• block 사이의 atomic 은 동기화가 아닌 update 만 — atomicAdd 가 race-free 하게 update 를 보장하지만, 순서를 보장하진 않는다. 이게 L009 Reductions 의 핵심 함정.

강의에서 Thomas 가 던지는 두 번째 정량. 한 cycle 안에 register 를 읽는 비용과 HBM 을 읽는 비용은 100~200배 차이다. on-chip 위계를 register → shared/L1 → L2 → HBM 으로 깔고 각 단계의 대략적 latency 와 bandwidth 를 외워둔다.

강의의 가장 중요한 한 그림. arithmetic intensity = FLOPs / byte. 같은 알고리즘이라도 이 값이 작으면 (memory-bound) bandwidth 가 천장이고, 크면 (compute-bound) FLOPs 가 천장이다. 두 천장이 만나는 “ridge point” 가 GPU 마다 정해져 있다 — A100 의 경우 약 13 FLOP/byte (FP32) 또는 ~140 FLOP/byte (Tensor Core BF16).

roofline 위에서 자주 보이는 패턴들.

• elementwise op (add, relu, sigmoid) 는 거의 항상 far left — bandwidth 가 천장이다. fusion 의 가장 큰 표적 (§08).
• matmul 은 어느 사이즈에서는 ridge 를 넘어 compute-bound. 하지만 tall-skinny matmul (예: M=2, K=1024, N=8192) 은 reuse 가 적어 memory-bound 로 떨어진다.
• attention 의 score 계산 (Q·Kᵀ) 은 큰 matmul 처럼 보이지만, softmax 와 V·attn 까지 다 묶어 계산할 때 HBM 왕복이 dominant. L012 FlashAttention 의 본론.
• 같은 모델의 같은 레이어가 training 시 bound 가 다르다 — backward 의 reduction step 이 종종 memory-bound. 한 화살표가 아니다.

occupancy = SM 의 active warp 수 / max warp 수. 이 비율이 100% 면 “SM 이 가득 찼다”. 그런데 이 비율을 결정하는 건 세 자원 중 가장 빡빡한 자리다.

강의에서 Thomas 가 깐 직관 — “higher occupancy 가 항상 좋은 건 아니다.” 이게 흔한 오해다. occupancy 는 memory latency 를 hide 하는 도구다. 한 warp 가 HBM 을 기다리는 동안 다른 warp 가 일을 하면 latency 가 hide 된다. 그래서 memory-bound 커널에서는 occupancy 가 중요하지만, 이미 compute-bound 면 occupancy 가 80% → 60% 떨어져도 거의 차이 없을 수 있다.

강의의 작은 실험. __global__ void empty() {} 같은 빈 커널을 launch 하면 — host → device 의 launch dispatch 만큼의 시간이 걸린다. 이게 “fixed cost per kernel launch”. 대략 5~20 μs (system, driver, GPU 종류에 따라).

이 수치가 의미 있는 자리는 — 커널이 작을수록이다.

마지막 줄이 L006 Optimizing Optimizers 의 핵심 동기다 — 10000 개의 작은 tensor 에 같은 op 를 따로 launch 하면 launch overhead 만 100 ms. foreach / multi_tensor_apply 가 이걸 한 launch 로 합쳐서 거의 0 으로 줄인다.

강의에서 Thomas 가 가장 명료하게 한 줄로 정리한 fusion 의 동기. “PyTorch eager 는 op 마다 input 을 read, 결과를 write 한다. 두 op 를 합치면 중간을 register 에 두고 한 번만 write.”

fusion 의 자동화는 두 자리에서 일어난다.

• torch.compile — Inductor 가 elementwise chain 을 자동 fuse. L001 §07 에서 TORCH_LOGS=output_code 로 검증.
• 커널 직접 짜기 — Triton/CUDA 로 여러 op 를 한 커널에 합친다. L018 Fusing Kernels 의 본론.

강의에서 Thomas 가 짠 GELU 커널이 좋은 예 — 0.5f * x * (1.0f + tanhf(sqrt(2/pi) * (x + 0.044715 * x³))) 안에 곱셈/덧셈/tanh 가 다 들어 있다. 이걸 PyTorch eager 처럼 풀면 5~6개의 elementwise launch + HBM 왕복이 되지만, 한 커널 안에서는 한 번의 read 와 한 번의 write 로 끝난다.

강의에서 Thomas 가 자기 디버깅 일화로 깐 가장 인상적인 자리. 옛날 Torch 에서 새 PyTorch 로 batchnorm 커널을 옮기다가 indexing 을 무심코 int64 로 했더니 새 커널이 더 느려졌다. 원인 추적에 한참 걸렸고 — register 사용량이 늘어나서 spill 이 생겼다는 게 진단이었다.

이 일화의 메시지가 두 줄로 압축된다.

1. indexing 의 dtype 도 성능 결정 요소. int 면 충분한 자리에 long 을 쓰지 마라. CUDA 의 blockIdx.x 등은 unsigned int.
2. “느려진 커널” 의 원인이 코드의 한 글자에 있을 수 있다 — 그래서 NCU 의 register/local memory 카운터를 항상 본다.

강의의 마지막 메시지가 이 일화에서 자연스럽게 따라 나온다 — “성능 일은 추측이 아니라 측정이다.” 그리고 측정하려면 roofline · occupancy · register pressure 의 단어들이 손에 잡혀 있어야 한다. 그게 이 강의의 모든 것.

PMPP Ch.4–5 의 정량 모델을 다시 펼치기 전에 머릿속에 남아 있어야 하는 사실들과 자료.

손에 새기기 — 실습 시퀀스

1. 자기 GPU 의 roofline 그리기 — peak FLOPs (FP32, FP16, Tensor Core), peak HBM bandwidth 를 spec 에서 가져와 ridge point 계산. matplotlib 으로 직접.
2. 빈 커널 시간 측정 — __global__ void empty(){} 을 launch 하고 cudaEvent 로 시간 측정. 자기 환경의 fixed cost 가 5 μs 인지 20 μs 인지 확정.
3. elementwise vs matmul 의 roofline 위치 — x + y 와 matmul(M=1024, N=1024, K=1024) 의 시간을 측정해 arithmetic intensity 와 함께 roofline 위에 점 찍기.
4. register sweep — 같은 커널을 일부러 변수 많이 만들어서 (nvcc -Xptxas=-v 로 register count 확인) occupancy 가 떨어지는 지점을 찾는다.
5. int32 vs int64 indexing 비교 — Thomas 의 batchnorm 일화 재현. 같은 reduce 또는 stencil 커널을 두 dtype 으로 짜고 시간 비교.
6. fusion 효과 측정 — relu(x) * 2 를 PyTorch eager 와 torch.compile 로 각각 돌려 시간 비교. TORCH_LOGS=output_code 로 fused 커널을 직접 본다.
7. NCU 로 occupancy 보기 — ncu --set full python script.py 의 output 에서 “Achieved Occupancy” 와 “Theoretical Occupancy” 의 차이를 한 번 본다.
8. 한 페이지 plan — 자기가 다루는 모델의 한 layer 를 골라 “이 layer 가 roofline 위에서 어디쯤 있는지, 어떤 천장을 깎으면 가장 효과적인지” 를 한 페이지로 정리.

L004 의 단어들 (occupancy, scoreboard stall, arithmetic intensity, fusion) 이 거의 모든 후속 강의의 기본 어휘다.

강의의 framing 안에서 의도적으로 비워둔 자리들과, 자기 GPU 위에서 측정해야 손에 박히는 사실들.

• Tensor Core 의 roofline 위치 — 강의는 FP32 unit 만 본다. Tensor Core 가 들어가면 ridge point 가 ~10x 오른쪽으로 이동. L023 CUTLASS 강의에서.
• arithmetic intensity 의 정확한 정의 — “FLOP / byte” 의 byte 가 무엇을 셀까? HBM read 만? L2 hit 까지 포함? 강의에서는 흐릿. NVIDIA 공식 정의는 “DRAM bytes” — 즉 HBM 만.
• occupancy 와 latency hiding 의 정량 — 100% occupancy 가 75% 보다 얼마나 빠른가? bound 에 따라 다름. 강의에서 정량적인 비교는 안 함.
• register pressure 의 정확한 측정 — nvcc -Xptxas=-v 가 보여주는 “registers per thread” 와 NCU 의 “registers per thread” 가 일치하는가? compiler 의 inline 결정 이후의 실제 사용량인가, 정의된 값인가?
• L2 cache 의 영향 — 강의는 L1/shared 와 HBM 만 깐다. L2 의 hit 률이 실제 bandwidth 에 어떻게 기여하는지는 빠진다. L008 와 NCU 의 “L2 cache throughput” 카운터에서.
• Hopper (H100) 의 SM 변화 — 강의는 Ampere 기준. H100 은 SM 당 더 많은 register, 더 큰 L2, async copy 의 등장 (cp.async). 같은 정량 모델이지만 숫자가 다르다.
• GELU 커널의 직접 측정 — Thomas 가 강의에서 GELU 커널을 짜고 PyTorch eager 와 비교했지만, 측정 결과를 정확히 재구성하지 못했다. (확인 필요) notebook 에서 직접 돌려 시간 차이 확인.