gpumode · 강의 아카이브
《GPU Mode》 L064 2025 · Multi-GPU High priority transcript · failed

Multi-GPU Programming

한 GPU 의 한계를 넘는 순간 — P2P · NVLink · NCCL 의 산수가 든다. Markus Hrywniak (NVIDIA) 의 강의는 multi-GPU 의 통신 모델을 baseline 부터 깐다 — peer-to-peer copy, stream + event 동기화, NCCL collectives, PyTorch DDP/FSDP 위 패턴, 그리고 흔한 함정. 이 노트는 그 sequence 를 한 노드 안multi-node 의 차이까지 정리한다.

P2P · NVLink NCCL collectives stream + event DDP · FSDP all-reduce / all-gather SHARP / NVSwitch topology-aware multi-node
M
Speaker
Markus Hrywniak
NVIDIA · Solution architect / DevTech
강의 번호
L064
스피커
Markus Hrywniak
학습 우선순위
High · 정독
자료 상태
transcript 없음 · NVIDIA 자료
§ 01강의가 풀려는 문제· when one GPU is not enough

한 GPU 가 꽉 찼을 때 — 두 GPU 가 한 GPU 처럼 안 보인다

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 / 코드는 영상 검증 필요.

강의의 frame

multi-GPU 의 핵심 어려움은 3가지 동기화. (1) 같은 GPU 의 두 stream 사이. (2) 같은 노드의 두 GPU 사이. (3) 다른 노드의 두 GPU 사이. 각 단계가 다른 latency / bandwidth / failure mode 를 가짐. PyTorch 의 추상이 이 셋을 다 같은 모양으로 보여주지만 — 성능을 짜내려면 차이를 다시 알아야 한다.

“multi-GPU 프로그래밍의 90% 는 통신 그 자체가 아니라 — 통신과 compute 를 어떻게 동시에 돌리느냐의 문제다.”Markus Hrywniak / 확인 필요
§ 02멀티 GPU 의 통신 모델· memory · bandwidth · latency

각 GPU 가 자기 HBM 을 갖는다 — 그래서 “데이터를 옮긴다” 가 명시적

single-GPU 프로그래밍의 가장 큰 추상 — host/device 의 두 메모리. multi-GPU 에서는 — 각 GPU 가 자기 HBM. 한 GPU 의 데이터를 다른 GPU 가 읽으려면 명시적 transfer. 이게 모든 distributed training / inference 의 cost 를 결정한다.

FIG · 통신 hierarchy 의 bandwidthH100 8-GPU 노드 + multi-node
SM ↔ HBM (intra-GPU)
~3.3 TB/s
NVLink (intra-node, P2P)
~900 GB/s
NVSwitch (intra-node, all-to-all)
~900 GB/s
PCIe Gen5 (host, intra-node)
~64 GB/s
InfiniBand NDR (inter-node)
~50 GB/s
Ethernet 100Gbps
~12 GB/s
HBM (intra-GPU) 은 NVLink (intra-node) 의 ~4배. NVLink 는 InfiniBand (inter-node) 의 ~18배. 한 단계 내려갈 때마다 BW 가 한 자리 수 떨어진다. 이게 모든 multi-GPU 디자인 결정의 기본 산수.

이 hierarchy 가 distributed 모델 학습의 layout 을 결정한다.

  • 같은 노드 안: NVLink. all-reduce, all-gather 같은 collective 가 거의 무료.
  • 다른 노드: InfiniBand 또는 더 느림. 통신을 compute 와 overlap 시키지 않으면 utilization 폭락.
  • cross-cluster (datacenter): ethernet. checkpoint 동기화 정도. training 중 통신은 안 함.

그래서 — TP (tensor parallel) 는 NVLink 가 있는 같은 노드 안. PP (pipeline parallel) 와 DP (data parallel) 는 노드 사이까지 OK. parallel 형태 선택이 곧 이 hierarchy 의 직접 반영.

NVLink 의 진화

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% 수준. 산수할 때 이 격차를 고려.

§ 03P2P / NVLink / NCCL· layer 위치

같은 통신을 표현하는 3개의 layer — 어디서 어떤 추상을 쓰는가

multi-GPU 통신은 layered. 같은 “GPU A 의 데이터를 GPU B 로” 가 — 어떤 layer 의 API 를 쓰느냐에 따라 다른 모양. 강의의 핵심 frame 중 하나가 이 layer 의 분리.

L0 · CUDA P2P

cudaMemcpyPeerAsync · 가장 raw

GPU A 의 device pointer 를 GPU B 의 device pointer 로 직접 copy. NVLink 가 있으면 자동으로 NVLink 사용.

cudaMemcpyPeerAsync(dst, dstDev, src, srcDev, n, stream)

가장 작은 단위. 한 쌍의 GPU 사이만. collective 없음.

L1 · NVLink + Direct Access

cudaDeviceEnablePeerAccess

peer access 를 켜면 — GPU A 의 커널이 GPU B 의 메모리를 직접 dereference 가능. kernel 안에서 다른 GPU 의 데이터에 접근.

매우 강력하지만 위험 — 동기화 / consistency 사용자가 직접 보장.

L2 · NCCL collectives

all-reduce · broadcast 등

여러 GPU 의 통신 패턴을 collective 로 추상화. topology-aware — NVLink / NVSwitch / InfiniBand 자동 사용.

ncclAllReduce(send, recv, n, dtype, op, comm, stream)

distributed training 의 표준. PyTorch dist.all_reduce 의 backend.

L3 · PyTorch distributed

DDP · FSDP · ProcessGroup

NCCL 위의 PyTorch 추상. autograd 와 통합되어 backward 의 gradient 동기화 자동.

대부분의 사용자가 만지는 자리. 그러나 디버깅 / 최적화 시 한 단계 아래로 내려가야 함.

언제 어디서

일반적인 ML training — L3 (PyTorch). 새로운 통신 패턴 (예: disaggregated KV transfer) — L2 (NCCL primitive). 매우 작은 latency 를 짜낼 때 (custom collective) — L1 또는 L0. 대부분의 PyTorch 사용자는 평생 L3 만 보면 되지만 — 한계까지 짜낼 때 한 단계씩 내려간다.

§ 04stream + event 동기화· async correctness

multi-GPU 의 동기화 — single-GPU 의 stream 모델을 multi-GPU 로 확장

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.

FIG · 두 GPU 의 stream + eventcompute / comm overlap
GPU 0 · compute
layer 0
layer 1
layer 2
GPU 0 · comm
send out0 → GPU1
send out1 → GPU1
GPU 1 · compute
— 대기 (e0)
layer 1
layer 2
comm 이 별도 stream — compute 와 겹친다. layer 1 의 send 는 layer 1 의 compute 이 끝날 때 시작 (event), 그 동안 layer 2 의 compute 가 도는 식.
기본 규칙

(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 에서는 드물게.

§ 05shared memory across GPUs· UVA · IPC

UVA — 하나의 가상 주소 공간이 여러 GPU 메모리를 다 본다

CUDA 의 UVA (Unified Virtual Addressing) 가 깐 추상 — 같은 process 안의 모든 GPU 메모리 + host 메모리가 한 가상 주소 공간. cudaMemcpyAsync 가 src/dst 의 device 를 자동 추론. cudaIPC 가 다른 process 의 GPU 메모리를 우리 process 가 보게 해줌 — vLLM 같은 multi-process 시스템의 기반.

vLLM 의 사례

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 시스템에서 자주 등장. 한 단계 더 들어가면 직접 만지게 되는 자리.

§ 06PyTorch 위 패턴· DDP · FSDP · TP/PP

같은 multi-GPU hardware 위 4가지 parallel 형태

FIG · 8-GPU 노드의 4가지 parallel 배치같은 hw, 다른 schedule
DP (8-way)
G0
G1
G2
G3
G4
G5
G6
G7
TP (8-way)
T0
T1
T2
T3
T4
T5
T6
T7
PP (8-stage)
P0
P1
P2
P3
P4
P5
P6
P7
2DP × 4TP
D0T0
D0T1
D0T2
D0T3
D1T0
D1T1
D1T2
D1T3
같은 8 GPU 가 — 모델 크기 / batch / latency 요구에 따라 4가지 다른 layout. 각 layout 이 다른 collective 패턴 (DP=allreduce gradient, TP=allreduce activation, PP=p2p activation). 큰 모델은 보통 hybrid (D × T × P) 로 묶음.
DDP — Data Parallel

같은 모델 N copy

각 GPU 가 모델 전체 + 다른 mini-batch. backward 후 gradient all-reduce. PyTorch DistributedDataParallel. 가장 단순.

한계 — 모델이 한 GPU 에 안 들어가면 사용 불가.

FSDP — Fully Sharded

parameter / gradient / optimizer shard

모델 weight 를 N 등분. 매 layer 마다 all-gather (forward) + reduce-scatter (backward). DeepSpeed ZeRO-3 / PyTorch FSDP.

거의 모든 큰 모델 학습의 표준.

TP — Tensor Parallel

matmul split

한 GEMM 의 column / row 를 N 등분. forward / backward 마다 all-reduce. NVLink 가 있는 노드 안에서만 효율적.

큰 모델 inference / training 의 흔한 첫 단계.

PP — Pipeline Parallel

layer split

모델의 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 패턴.

§ 07흔한 함정· deadlock · blocking · stale data

multi-GPU 디버깅의 표준 함정 list

“multi-GPU 의 버그는 single-GPU 의 버그보다 — 한 단계 더 visible 하지 않다. timeline 도구가 거의 항상 답.”학습 노트
§ 08성능 디버깅· nsys · NCCL_DEBUG

multi-GPU timeline 을 어떻게 보는가

도구 첫 번째 — 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
§ 09multi-node 와의 차이· InfiniBand · GPUDirect RDMA

한 노드를 넘는 순간 — bandwidth 가 한 자리수 떨어진다

실전 layout 결정

(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. 큰 모델 학습의 표준.

§ 10기억할 메모와 코드· key takeaways

다시 열었을 때 빠르게 잡혀야 할 것

통신 hierarchy
HBM > NVLink > PCIe > InfiniBand > ethernet. 각 단계 ~10× 차이.
layered API
P2P → peer access → NCCL → PyTorch DDP. 한계까지는 한 단계씩 내려간다.
stream + event
cross-GPU 동기화. host 의 sleep 절대 금지. event.wait() 가 GPU 측 대기.
UVA · cudaIPC
한 process 안 / 다른 process 사이의 GPU 메모리 공유. vLLM 류 시스템 기반.
DP / TP / PP / FSDP
4가지 parallel 형태. 보통 hybrid (DP × TP × PP) 로 묶음.
TP=노드 안만
NVLink 의 18× 우위. cross-node TP 는 거의 항상 잘못된 결정.
함정 list
collective mismatch deadlock / topology mismatch / default stream 직렬화 / cudaIPC dangling.
디버깅 도구
NCCL_DEBUG=INFO/TRACE, nsys timeline, nccl-tests baseline.
NCCL docs NCCL User Guide
PyTorch distributed torch.distributed · DDP / FSDP / PG
Megatron-LM github.com/NVIDIA/Megatron-LM · TP/PP/DP hybrid 의 reference

손에 새기기 — 실습 시퀀스

  1. nvidia-smi topo -m — 자기 시스템의 GPU topology 확인. NVLink 연결, NUMA 도메인, PCIe path.
  2. nccl-tests all-reduce — 자기 시스템의 baseline BW. peak 의 80%+ 도달 못 하면 topology 문제.
  3. P2P memcpy 직접cudaMemcpyPeerAsync 로 GPU 0 → GPU 1 transfer. nsys 로 NVLink 사용 확인.
  4. DDP toy training — 4-GPU 로 ResNet 학습. NCCL_DEBUG=INFO 로 collective trace.
  5. FSDP 70B 시도 — 8-GPU 로 70B Llama. param shard / gather pattern 의 timeline 확인.
  6. TP × DP hybrid — 8 GPU 를 4 TP × 2 DP. Megatron-LM 또는 PyTorch native 로.
  7. cross-node 시도 — 두 노드 사이 InfiniBand. 같은 collective 가 노드 안 vs 노드 사이 BW 비교.
  8. collective mismatch deadlock 재현 — rank 별 다른 shape 의 all-reduce 호출. 의도적으로 deadlock. NCCL 의 timeout / hang 동작 확인.
§ 11다른 강의로· connections

이 강의의 frame 이 다른 강의에서 어떻게 다시 등장하는지

§ 12열린 질문· open questions

transcript 가 비어 있어 직접 검증해야 할 것들

검증 메모

본 노트의 모든 수치 (BW 절대값, 18× 차이 등) 는 H100 / NVLink4 / InfiniBand NDR 의 generic 산수. 자기 hw 세대 / topology 에 따라 절대값 변동. 항상 nccl-tests 와 nvidia-smi topo 로 확인 후 의사결정.

← Lecture 063 Search-Based Compilers — single-GPU schedule 의 다른 axis Lecture 065 → Neighborhood Attention — multi-GPU 보다 single-kernel sparsity