한 GPU 의 한계를 넘는 순간 — P2P · NVLink · NCCL 의 산수가 든다. Markus Hrywniak (NVIDIA) 의 강의는 multi-GPU 의 통신 모델을 baseline 부터 깐다 — peer-to-peer copy, stream + event 동기화, NCCL collectives, PyTorch DDP/FSDP 위 패턴, 그리고 흔한 함정. 이 노트는 그 sequence 를 한 노드 안 과 multi-node 의 차이까지 정리한다.
PyTorch 의 자동 multi-GPU 추상 (DDP, FSDP) 이 너무 잘 도는 시대 — 그래서 사람들이 multi-GPU 프로그래밍의 기본 모델을 잊는다. 이 강의는 “PyTorch 가 우리 대신 무엇을 하고 있는가” 를 한 단계 아래에서 본다 — CUDA P2P, stream + event, NCCL collectives 의 layer 가 어떻게 쌓이는지.
강의 transcript 가 비어 있으므로 본 노트는 NVIDIA 의 multi-GPU programming 자료 (CUDA programming guide ch. 6, NCCL docs, GTC 강의), PyTorch distributed 자료, 그리고 일반적 GPU systems 자료 위에서 재구성한다. Markus 가 강의에서 강조한 정확한 demo / 코드는 영상 검증 필요.
multi-GPU 의 핵심 어려움은 3가지 동기화. (1) 같은 GPU 의 두 stream 사이. (2) 같은 노드의 두 GPU 사이. (3) 다른 노드의 두 GPU 사이. 각 단계가 다른 latency / bandwidth / failure mode 를 가짐. PyTorch 의 추상이 이 셋을 다 같은 모양으로 보여주지만 — 성능을 짜내려면 차이를 다시 알아야 한다.
single-GPU 프로그래밍의 가장 큰 추상 — host/device 의 두 메모리. multi-GPU 에서는 — 각 GPU 가 자기 HBM. 한 GPU 의 데이터를 다른 GPU 가 읽으려면 명시적 transfer. 이게 모든 distributed training / inference 의 cost 를 결정한다.
이 hierarchy 가 distributed 모델 학습의 layout 을 결정한다.
그래서 — TP (tensor parallel) 는 NVLink 가 있는 같은 노드 안. PP (pipeline parallel) 와 DP (data parallel) 는 노드 사이까지 OK. parallel 형태 선택이 곧 이 hierarchy 의 직접 반영.
V100 (NVLink2 = 300 GB/s) → A100 (NVLink3 = 600 GB/s) → H100 (NVLink4 = 900 GB/s) → B200 (NVLink5 = 1800 GB/s). 한 세대마다 ~2배. NVSwitch 가 추가되면 8 GPU 가 모두 동시에 peer 와 full-BW 통신 가능. “같은 노드 안 = single-GPU 처럼 보이는 통신” 의 hardware 기반.
주의 — NVLink 의 “900 GB/s” 는 peer 한 쌍의 양방향 합. 단방향이면 450 GB/s. 그리고 noise / overhead 로 실제 NCCL 측정값은 80~85% 수준. 산수할 때 이 격차를 고려.
multi-GPU 통신은 layered. 같은 “GPU A 의 데이터를 GPU B 로” 가 — 어떤 layer 의 API 를 쓰느냐에 따라 다른 모양. 강의의 핵심 frame 중 하나가 이 layer 의 분리.
GPU A 의 device pointer 를 GPU B 의 device pointer 로 직접 copy. NVLink 가 있으면 자동으로 NVLink 사용.
cudaMemcpyPeerAsync(dst, dstDev, src, srcDev, n, stream)
가장 작은 단위. 한 쌍의 GPU 사이만. collective 없음.
peer access 를 켜면 — GPU A 의 커널이 GPU B 의 메모리를 직접 dereference 가능. kernel 안에서 다른 GPU 의 데이터에 접근.
매우 강력하지만 위험 — 동기화 / consistency 사용자가 직접 보장.
여러 GPU 의 통신 패턴을 collective 로 추상화. topology-aware — NVLink / NVSwitch / InfiniBand 자동 사용.
ncclAllReduce(send, recv, n, dtype, op, comm, stream)
distributed training 의 표준. PyTorch dist.all_reduce 의 backend.
NCCL 위의 PyTorch 추상. autograd 와 통합되어 backward 의 gradient 동기화 자동.
대부분의 사용자가 만지는 자리. 그러나 디버깅 / 최적화 시 한 단계 아래로 내려가야 함.
일반적인 ML training — L3 (PyTorch). 새로운 통신 패턴 (예: disaggregated KV transfer) — L2 (NCCL primitive). 매우 작은 latency 를 짜낼 때 (custom collective) — L1 또는 L0. 대부분의 PyTorch 사용자는 평생 L3 만 보면 되지만 — 한계까지 짜낼 때 한 단계씩 내려간다.
CUDA stream 은 order constraint 의 단위. 같은 stream 안의 일은 순서대로, 다른 stream 끼리는 병렬. multi-GPU 에서는 각 GPU 가 자기 stream 을 갖는다. cross-GPU 동기화를 위해 event 를 쓴다.
# GPU 0 의 결과를 GPU 1 이 기다린다
with torch.cuda.device(0):
out0 = layer0(x)
e0 = torch.cuda.Event()
e0.record() # GPU0 stream 에 마커
with torch.cuda.device(1):
e0.wait() # GPU1 stream 이 e0 끝까지 대기
# 이 시점부터 out0 는 GPU1 위에 visible
out1 = layer1(out0.to(1))
중요한 점 — e0.wait() 은 host 가 block 안 된다. GPU1 의 stream 이 GPU0 의 e0 가 끝날 때까지 대기. host 는 다음 일 진행. 이 비동기성이 multi-GPU 의 throughput 의 source.
(1) cross-GPU dependency 는 event 로. host 의 sleep / busy-wait 절대 금지. (2) comm 은 별도 stream. default stream 에 두면 compute 와 직렬화. (3) stream 마다 priority 가능 — 작은 latency-critical comm 은 high-priority stream 에. (4) cudaStreamSynchronize 는 디버깅용 — production 에서는 드물게.
CUDA 의 UVA (Unified Virtual Addressing) 가 깐 추상 — 같은 process 안의 모든 GPU 메모리 + host 메모리가 한 가상 주소 공간. cudaMemcpyAsync 가 src/dst 의 device 를 자동 추론. cudaIPC 가 다른 process 의 GPU 메모리를 우리 process 가 보게 해줌 — vLLM 같은 multi-process 시스템의 기반.
cudaPointerGetAttributes(ptr) 로 어디 주소인지 확인 가능. peer access 가 켜진 GPU 끼리는 서로의 주소를 직접 dereference.cudaIpcGetMemHandle / cudaIpcOpenMemHandle. socket 으로 주고받음.vLLM 의 multi-process tensor-parallel 구조 — 각 GPU 가 별도 process. KV cache 가 cudaIPC handle 로 공유 (또는 같은 process tree 안에서 fork). 그래서 새 prefill 요청이 들어왔을 때 KV slot 을 다른 process 가 직접 write 가능. process 분리 + GPU 메모리 공유 의 흔한 패턴.
UVA / IPC 는 — distributed training 의 high-level NCCL 위로 잘 안 보이지만, KV transfer / disaggregation / inference 시스템에서 자주 등장. 한 단계 더 들어가면 직접 만지게 되는 자리.
각 GPU 가 모델 전체 + 다른 mini-batch. backward 후 gradient all-reduce. PyTorch DistributedDataParallel. 가장 단순.
한계 — 모델이 한 GPU 에 안 들어가면 사용 불가.
모델 weight 를 N 등분. 매 layer 마다 all-gather (forward) + reduce-scatter (backward). DeepSpeed ZeRO-3 / PyTorch FSDP.
거의 모든 큰 모델 학습의 표준.
한 GEMM 의 column / row 를 N 등분. forward / backward 마다 all-reduce. NVLink 가 있는 노드 안에서만 효율적.
큰 모델 inference / training 의 흔한 첫 단계.
모델의 layer 를 N stage 로. 각 GPU 가 자기 stage 만. p2p send / recv. micro-batch 로 bubble 줄임.
cross-node 까지 OK — comm 이 작음. 큰 모델의 절대 필요.
실제 production 모델 학습은 거의 항상 — DP × TP × PP 의 hybrid. 예를 들어 H100 256-GPU 클러스터에서 70B Llama 학습 — DP=8 × TP=8 × PP=4. 각 axis 가 다른 collective 패턴.
all_reduce(A), rank 1 이 all_reduce(B). 두 텐서 shape 다름. NCCL 이 영원히 대기. 해결 — 모든 rank 가 같은 collective 를 같은 shape 로 호출.nvidia-smi topo -m 으로 topology 확인.PYTORCH_CUDA_ALLOC_CONF 튜닝).도구 첫 번째 — NCCL 환경변수:
NCCL_DEBUG=INFO — 통신 algorithm / topology 선택 출력. ring vs tree, 어떤 NVLink 사용. 첫 sanity check.NCCL_DEBUG=TRACE — 모든 collective 의 시작/끝 timestamp. very verbose.NCCL_TOPO_DUMP_FILE — topology 의 JSON dump. 자기 클러스터의 hw 구성을 확인.NCCL_ALGO=ring|tree, NCCL_PROTO=simple|LL|LL128 — 알고리즘 강제.도구 두 번째 — nsys (Nsight Systems). multi-GPU timeline 을 한 화면에. GPU lane 별로 compute / memcpy / NCCL 의 각 연산을 segment 로. NVTX range 를 자기 코드에 박으면 더 정확히 매핑.
(1) nsys 로 trace — 한 step 의 timeline. (2) collective 의 시작 / 끝 정렬 — 모든 rank 가 같은 시점에 collective 를 entered? 아니면 한 rank 가 늦어서 다른 rank 가 기다림? (3) compute idle gap — 어디서 어느 GPU 가 놀고 있는지. (4) NCCL bandwidth 측정 — 측정한 전송 시간 / 데이터 크기 = 실효 BW. peak 의 80~85% 가 정상.
NCCL 의 표준 micro-benchmark — nccl-tests. 같은 cluster 에서 all-reduce / all-gather 의 BW 를 size 별로 측정. 자기 시스템의 baseline 을 잡을 때 첫 step.
# NCCL 의 표준 micro-benchmark
$ NCCL_DEBUG=INFO mpirun -np 8 ./build/all_reduce_perf -b 8 -e 256M -f 2
# 출력 (대략)
size count busbw (peak) time(us)
8 2 0.00 12.5
1024K 256K 145.34 720.8
16384K 4194K 210.50 19200 # peak 의 87%
268435K 67108K 218.34 230000
(1) TP 는 한 노드 안. cross-node TP 는 거의 항상 잘못된 결정 — NVLink 의 18× 우위. (2) DP 는 cross-node OK. gradient all-reduce 는 latency-tolerant. (3) PP 는 cross-node 도 OK. p2p send/recv 가 작음. (4) hybrid: 한 노드 안 TP=8, 노드 사이 DP+PP. 큰 모델 학습의 표준.
cudaMemcpyPeerAsync 로 GPU 0 → GPU 1 transfer. nsys 로 NVLink 사용 확인.NCCL_DEBUG=INFO 로 collective trace.본 노트의 모든 수치 (BW 절대값, 18× 차이 등) 는 H100 / NVLink4 / InfiniBand NDR 의 generic 산수. 자기 hw 세대 / topology 에 따라 절대값 변동. 항상 nccl-tests 와 nvidia-smi topo 로 확인 후 의사결정.