분산 통신 & 병렬화 단권화

1 Collective 정의

• point-to-point (SEND/RECV) 의 상대개념
• 의미론: (1) input set, (2) output set, (3) reduction op (sum/min/max/prod)
• communicator 단위로 정의 — rank 0..N-1

2 AllReduce ★ 가장 중요 y[i] = Σ x_r[i]

• DP gradient sync · TP row-parallel 뒤 합산
• = ReduceScatter ∘ AllGather (분해 가능)

3 Reduce 단일 root

실무: 로그·통계 집계. 학습에는 거의 안 씀 (AllReduce 대체).

4 Broadcast ↔ Reduce 쌍

• initial weight 배포 · random seed 동기화
• 통신량: (N−1)·S · 이론 하한 step = ⌈log₂N⌉ (tree)

5 AllGather ★

• ZeRO-3 / FSDP param unshard  ↗ V17 §3
• TP column-parallel 입력 broadcast 대체
• ring: byte/GPU = (N−1)/N · S

6 ReduceScatter ★

• ZeRO-2 grad sharding · FSDP backward
• AllReduce = AllGather ∘ ReduceScatter (분해 등식)

7 AllToAll ★ shuffle

• EP (MoE) token routing · SP (Ulysses) head↔seq 전치
• 통신량: (N−1)/N · S · ring / Bruck / pairwise 알고리즘

8 Collective 의미론 매트릭스

S = total message size in bytes. N = rank 수. NCCL: Scatter/Gather는 AllGather/ReduceScatter 우위라 거의 안 씀.

9 분해 등식 (!) AR=AG∘RS

1 Ring AllReduce ★ bandwidth-optimal RS→AG

1. Phase A — ReduceScatter: N−1 step, 각 step 마다 chunk 1개 전송·수신 후 local reduce 누적. 끝에 각 rank는 chunk 1개를 "완성 형태"로 보유.
2. Phase B — AllGather: N−1 step, 완성된 chunk를 링으로 회전. 끝에 모두 동일 전체 결과 S 보유.

2 Ring 단계별 도식 (N=4)

C=S/4 ; rank r 가 chunk r 의 owner.

Phase A (ReduceScatter, step 0..N-2):
 step0:  r → r+1  send chunk (r)     recv chunk (r-1)    local += recv
 step1:  r → r+1  send chunk (r-1)   recv chunk (r-2)    local += recv
 step2:  r → r+1  send chunk (r-2)   recv chunk (r-3)    local += recv
 ─ after N-1=3 step:  r0 보유 chunk 1(완성) , r1→2 , r2→3 , r3→0

Phase B (AllGather, step 0..N-2):
 step0:  r → r+1  send 완성 chunk    recv 상류 완성 chunk
 step1:  r → r+1  send 방금 받은 것  recv 상류의 상류 것
 step2:  r → r+1  …
 ─ after N-1=3 step:  모든 rank = [c0,c1,c2,c3] full S

총 step = 2(N−1). 각 step 전송량 = C = S/N. rank당 총 전송 = 2(N−1)·S/N → BW-optimal.

3 Ring 복잡도

• S 클수록 ring 승 (β-term 우세)
• N·α 커서 작은 S·큰 N 에선 불리 — tree로 전환

4 Tree AllReduce latency-optimal small

small S / large N 에서 ring 보다 좋음.

5 Double Binary Tree ★ NCCL default 종종

• link 양방향 full duplex 활용 → intra-tree BW = full
• NCCL: 2-node+ 에서 자주 선택됨
• step ≈ 2⌈log N⌉ · bytes/GPU ≈ 2S/ log N 수준 (algo-specific)

6 Recursive Halving-Doubling

단 N 이 2의 거듭제곱일 때 깔끔. non-power-of-2 는 virtual rank 조정 필요.

7 알고리즘 비교 ★

bytes/GPU 는 양방향 합. NCCL 은 tuner로 자동 선택 · NCCL_ALGO 로 강제 가능.

8 하한 이론 (!) 2(N-1)/N · S

1 NCCL 계층 용어

2 Communicator 생성 흐름

1. ncclGetUniqueId (rank 0) → MPI/file 로 전파
2. 각 rank가 ncclCommInitRank(id, N, r)
3. 내부: topology detection · graph search · channel 빌드
4. 완료 시 comm handle = 통신 plan 전체를 담고 있음

초기화 비용 O(N) ~ 수백 ms. 재사용 필수.

3 Channel ★ 병렬 ring

• channel 수 = 기본 4~32 (topology·size에 따라 자동)
• NCCL_NCHANNELS · NCCL_MIN_NCHANNELS 로 조정
• 많을수록 small msg 에선 비싼 초기화 + SM 낭비

4 Topology Detection

• source: /sys/class/nvidia, nvidia-smi topo -m 내부 API, HWLOC
• NIC 는 PCIe root 기준으로 CPU affinity·GPU locality 판단
• 결과로 ring/tree graph 후보 집합 생성

5 Ring 구성 알고리즘

1. candidate Hamilton cycle 탐색 (GPU 노드를 한 번씩 방문)
2. edge 가중치 최소화: NVLink 선호 · PCIe 회피 · NIC traversal 최소
3. multi-node: intra-node local ring + inter-node stitch (single NIC per rail)
4. 결과: graph->inter/intra 경로 확정

6 Protocol 차이 ★

NCCL_PROTO=LL,LL128,Simple 조합 지정 가능.

7 Proxy Thread

• GPUDirect RDMA 가능 시 GPU 메모리 직접 전송 (proxy는 메타만) ↗ §16
• bottleneck 구간: proxy SM ↔ NIC PCIe lane

8 Kernel Launch 경로

host API (ncclAllReduce)
 └─ enqueue (graph+args)
    └─ kernel launch (stream)
       ├─ intra-node: P2P load/store NVLink
       └─ inter-node: FIFO write → proxy → NIC
                                        └─ NET send/recv
       └─ reduce (FMA in-kernel) + progress

collective 1회 = 1 CUDA kernel launch. CUDA Graph capture 가능.

9 Graph Mode & Capture

• NCCL 2.11+ CUDA Graph capture 지원 · launch overhead 제거
• decode path 에 유용  ↗ V16 §10
• NCCL_GRAPH_MIXING_SUPPORT=1

1 NVLink 세대별 BW

공식 스펙 (per-GPU aggregate, both directions). PCIe Gen5 x16 = 128 GB/s 와 비교 시 7× 이상.

2 PCIe vs NVLink ★

3 NVSwitch fabric

• H100 NVSwitch: 64 port NVLink4, switch당 aggregate 3.2 TB/s
• 8-GPU DGX 내부: any-to-any non-blocking

4 DGX H100 8-GPU topology

              NVSwitch (x4)
       ╔══════════════════════════╗
       ║   full mesh crossbar     ║
       ╚╤═╤═╤═╤═╤═╤═╤═╤══════════╝
        │ │ │ │ │ │ │ │
       G0 G1 G2 G3 G4 G5 G6 G7
        │                    │
       NIC0 ... (8 CX-7 NIC, 400G IB)
        │
       ── ConnectX-7 → 2 rail (rail-optimized)

각 GPU → 2 NIC (via PXN). NVLink4 900 GB/s bi-dir intra-node.

5 Multi-node: IB Fabric

• node 당 8 GPU + 8 CX-7 NIC (각 400 Gb/s = 50 GB/s)
• rail-optimized: 같은 rail 의 NIC 끼리 leaf switch 공유 → inter-node same-rail 저지연
• node당 NIC aggregate = 400 GB/s 수준 → NVLink 대비 1/2 이하

6 NVLink BW 이론 활용 ring

7 왜 full-mesh 가 필요한가

• TP + EP 조합: alltoall 빈발 → fabric 필수
• NVSwitch 없는 8-way 보드: ring 만 full BW, 나머지는 절반 이하

8 비(非) DGX 구성 warning

1 NVLS 개념 ★ switch-in reduce

• 기존 IB SHARP 의 NVLink fabric 판
• AllReduce / ReduceScatter / Broadcast 가속

2 기존 Ring 과의 차이

Ring (no NVLS):
  GPU0 → GPU1 → GPU2 → ... → GPU_{N-1} (N-1 step)
   각 GPU 가 부분합을 수행

NVLS (in-switch):
  GPU_r ──┐
  GPU_r' ─┼─► switch (ALU: sum) ──► all GPUs
  GPU_r'' ┘
   1 shot multicast-reduce

3 하드웨어 요건

• H100 SXM + NVSwitch 3 (DGX H100 / NVL72)
• CUDA 12.0+ · NCCL 2.17+
• mmap multicast object (CUmemMulticast) 지원
• Hopper PTX: multimem.ld_reduce, multimem.st

4 Bandwidth 이득 ★

• 실효 BW 2× 이득 (이론상)
• 실측: large msg 에서 1.5~1.9× 가속

5 Multicast Object

CUmemGenericAllocationHandle h;
CUmemAllocationProp prop = {...};
cuMulticastCreate(&mc, &mc_prop);
cuMulticastAddDevice(mc, dev);
cuMulticastBindAddr(mc, addr, size, 0);
// multimem.* PTX 로 접근

6 multimem PTX

7 NCCL 활성화

• NCCL_ALGO=NVLS · NCCL_NVLS_ENABLE=1
• comm init 시 자동 감지 — H100 + driver + nvidia-peermem OK 이면 활성
• 제한: intra-node 만 (single NVSwitch fabric)
• multi-node: CollNet SHARP (IB SHARP) 별도 → ↗ §16

8 용도별 효과

9 Trade-off 주의

1 문제 정의

2 Rail-optimized 란

Leaf switches: L0 L1 L2 ... L7 (rail 별 1개)
node A : G0→NIC0→L0 · G1→NIC1→L1 · ...
node B : G0→NIC0→L0 · G1→NIC1→L1 · ...
 같은 rank index 끼리는 같은 leaf → 1 hop
 다른 rank 끼리는 spine switch 경유 → 2+ hop

학습 시 AllReduce 는 ring 구성상 주로 same-rank pair → rail-matched 가 최적.

3 PXN 아이디어 ★ (!) local P2P then NET

4 PXN 흐름도

GPU_r (rail i, 목적지 rail j)
   │ ① NVLink P2P
   ▼
GPU_k (같은 node, rail j)
   │ ② RDMA WRITE via NIC_j
   ▼
NET (leaf-L_j 직행)
   ▼
GPU_r' (remote node, rail j)

비교 (no-PXN):
GPU_r → NIC_i → spine → L_j → GPU_r'  (2 hop)

5 효과 수치

6 NCCL 활성화

• NCCL_PXN_DISABLE=0 (기본값) — auto 활성
• NVLink P2P 가능 pair 간에만 동작
• ACS 이슈 없어야 함 (ACS = PCIe Access Control Services, 켜져있으면 P2P 차단)

7 NUMA & CPU affinity

• DGX: 2 NUMA · 4 GPU per NUMA
• PXN 이 cross-NUMA 만큼은 NVLink 로 우회

8 Rail 활용 체크리스트

1. GPU ↔ NIC 쌍 확인: nvidia-smi topo -m
2. IB 인터페이스 bind: NCCL_IB_HCA, NCCL_IB_GID_INDEX
3. rail 인덱스 일치: node 간 같은 rank → 같은 leaf
4. CPU affinity: proxy thread 를 GPU-local NUMA 에 pin
5. ACS off · IOMMU pass-through

9 Intra-node 변형

1 TP 정의 matmul 분할

• Megatron-LM (Shoeybi 2019) 가 표준
• TP group = 동일 layer 를 공유하는 rank 집합

2 Column-parallel ★

3 Row-parallel ★

4 Megatron MLP 패턴 ★ col→row→AR

MLP(x) = GeLU(x · W₁) · W₂

 W₁ : column-parallel (K×4H split N)
   y = GeLU(x · W₁ᵣ)     # output partition
 W₂ : row-parallel (4H×K split K)
   z = y · W₂ᵣ
   out = AllReduce(z)      # ← 1 AllReduce / MLP

activation (GeLU) 는 element-wise → partition 상태 유지 OK.

5 Attention 패턴

QKV proj   : column-parallel (head 차원 분할)
 each GPU → head subset 의 Q,K,V

attention  : local (heads 독립)

out proj   : row-parallel + AllReduce
 ← 1 AllReduce / attention block

head 수가 P 의 배수 여야 깔끔. GQA/MQA 에선 KV head 수 제약  ↗ V07 §6.

6 통신량 (layer 1회)

7 TP 크기 선택

8 왜 AllReduce 위치가 그런가

9 SP 와의 겹침 (!) SP=RS+AG

1 PP 정의 layer 분할

• micro-batch M 개 로 batch 쪼갬 → pipeline 채움
• 통신 = stage 경계 activation P2P · backward grad P2P

2 GPipe naive

stage↓  time →
 0 :  F1 F2 F3 F4 ..          B4 B3 B2 B1
 1 :     F1 F2 F3 F4 ..      B4 B3 B2 B1
 2 :        F1 F2 F3 F4 ..  B4 B3 B2 B1
 3 :           F1 F2 F3 F4 .B4 B3 B2 B1
 forward 전부 → backward 전부 (all-forward-all-backward)

• 단순 · 하지만 activation 전부 보관 → 메모리 ↑
• bubble ratio = (P−1)/M

3 1F1B ★ PipeDream / Megatron

stage↓  time →
 0: F1 F2 F3 F4 B1 F5 B2 F6 B3 F7 B4 ..
 1:    F1 F2 F3 F4 B1 F5 B2 F6 B3 F7 ..
 ...
   F 와 B 를 1 micro-batch 씩 교대

• memory: stage 당 보관 activation ≈ P (GPipe 는 M)
• bubble ratio = (P−1)/M (동일)
• throughput 유사, memory 우위

4 Interleaved 1F1B ★ Megatron virtual

5 Zero Bubble 최근

• backward 를 B (input grad) + W (weight grad) 로 쪼갬
• W 는 critical path 밖 → bubble 채움
• 조건 만족 시 bubble = 0 (수식상)

6 Bubble 공식 비교

Megatron 기본: v=2~4. throughput 개선 vs 통신 overhead trade.

7 통신 패턴

• stage 경계: SEND/RECV activation (forward) + grad (backward)
• P2P message — NCCL ncclSend/ncclRecv
• size: B · S · H per micro-batch per boundary
• NVLink intra-node / IB inter-node

8 PP 크기 선택 (!) M ≥ 4P

9 PP & 학습 loop

학습 loop 통합·FSDP 연동은 ↗ V17 §10 에서 상세. 본 권은 통신 패턴에만 초점.

1 Vanilla DP

2 학습 메모리 구성 Adam 기준

자세한 mixed precision 은 ↗ V09 §2 · 학습 loop 통합 ↗ V17 §2.

3 ZeRO-1 optimizer shard

4 ZeRO-2 + grad shard

5 ZeRO-3 / FSDP ★ + param shard

PyTorch FSDP = ZeRO-3 native · unit = FlatParameter (여러 param concat).

6 FSDP 한 layer 흐름 ★

layer forward:
  AllGather(param)        ← full W 복원 (N→1)
  y = layer(x)
  free param shard-only   ← shard 로 복귀

layer backward:
  AllGather(param)        ← 다시 필요
  dx, dw = layer.bwd(...)
  ReduceScatter(dw)       ← grad shard 로
  free param

• layer 마다 prefetch 로 AG overlap
• use_orig_params=True (FSDP-2) 로 inference 호환

7 Offload 변형

• ZeRO-Offload: CPU 로 opt state 내림
• ZeRO-Infinity: NVMe 까지 확장
• 제한: CPU-GPU BW (PCIe Gen4 ~32 GB/s) 가 bottleneck

8 비교 매트릭스

N = DP rank 수. 상세 수식 유도 ↗ V17 §3.

1 MoE 배경

• param ↑↑ (E·MLP_size) · FLOP 는 k/E 배
• Switch Transformer (k=1) · Mixtral (k=2, E=8) 등

2 EP 정의

3 핵심 통신 = AllToAll ★

MoE forward:
 x [B·S, H]                       각 rank 보유 token
   └─ router → assignment [B·S, k]
   └─ sort by expert
   └─ AllToAll (dispatch)  ← 1
   └─ expert.MLP (local)
   └─ AllToAll (combine)   ← 2
   └─ scale by gate · sum top-k

1 layer 당 2× AllToAll

4 Capacity Factor

• cf 높이면 drop ↓, 통신량 ↑
• cf 낮추면 GPU 활용도 ↑, 성능 손실 위험

5 AllToAll 크기

Mixtral 8x7B (k=2, E=8, H=4096): B·S = 8192 · 2·H=16KB/token · bytes ≈ 130 MB/rank/layer.

6 Token Routing 스키마

7 AllToAll 구현 알고리즘

• Ring: N−1 step, neighbor 전송. BW-optimal.
• Bruck: log N step, index reorder 필요. latency-optimal small msg.
• Pairwise exchange: N step. NVSwitch 하에서 단순.
• NCCL 내부 hierarchical: intra-node pairwise + inter-node ring

8 EP + TP 조합

• typical: EP=8, TP=1 (expert 단위로 분할)
• EP > 8 → inter-node AllToAll 병목

9 Load Balance Loss (!) uniform→no drop

1 SP 동기 TP 한계

• Megatron-SP: TP 그룹 내부에서 norm/dropout 을 seq 차원 shard
• Ulysses (DeepSpeed): 독립적 SP — seq 차원 shard, head 차원 AllToAll

2 Megatron-SP ★ TP 보조

기존 TP:
 [B,S,H] → AllReduce(TP) → [B,S,H] (전 rank 복제)

Megatron-SP:
 [B,S,H] → ReduceScatter(S축) → [B,S/P,H]
         → LayerNorm/Dropout (local)
         → AllGather(S축) → [B,S,H]
         → TP attn/MLP
AR = RS + AG 분해 → 추가 통신 없음, activation ÷P

3 Ulysses ★ 독립 SP

[B, S/P, H]
   │ AllToAll (S ↔ Head)
   ▼
[B, S, H/P]   ← head 차원 분할
 attention (per head 독립)
   │ AllToAll (Head ↔ S)
   ▼
[B, S/P, H]

4 Ulysses 통신량

5 SP & 활성 메모리

6 제약 head 수

7 Norm 과의 상호작용

• LayerNorm/RMSNorm 은 hidden 축 평균 → shard 상태에서도 local 계산 가능 (hidden 축 full)
• Megatron-SP: norm 앞에 RS (S축 shard) · 뒤에 AG (복원)
• Dropout mask = random → rank 간 seed 동기화 필수 (같은 shard 는 같은 mask)

8 TP + SP 통합 흐름

input [B,S/P,H]
 ├─ AllGather → [B,S,H]
 ├─ attention QKV (col-TP) → [B,S,H/P]
 ├─ attention (head 독립)
 ├─ out proj (row-TP) → ReduceScatter → [B,S/P,H]
 ├─ dropout (local, seed sync)
 ├─ layernorm (local)
 ├─ MLP W1 (col-TP) → AllGather → [B,S,H/P_mlp]
 ├─ GeLU (local)
 └─ MLP W2 (row-TP) → ReduceScatter → [B,S/P,H]

AR = RS+AG 로 분해해 SP 결합. 총 통신량은 TP-only 와 동일, activation 는 shard.

9 SP 크기 가이드

• 일반: SP = TP (같은 그룹 재사용)
• seq 매우 길면: Ulysses SP 독립 확장 · 혹은 CP 결합
• attention 커널 호환성 확인 (FlashAttn varlen)  ↗ V07 §9

1 CP 동기 long seq

• Q/K/V 를 seq 축으로 P 등분
• attention 은 K/V 를 ring 으로 순환 시켜 모든 Q 가 모든 K/V 를 만남

2 Ring Attention ★ 핵심

GPU_r 보유: Q_r, K_r, V_r  (seq shard r)

for t in 0..P-1:
  K_{r-t}, V_{r-t} 로 local attention (Q_r × K_{r-t}·V_{r-t})
  partial softmax + weighted sum 누적
  P2P SEND K,V → rank r+1 ; RECV from rank r-1

FlashAttn-style online softmax 로 수치 안정

P2P 전송과 attention 연산 overlap 가능 → BW 덜 드러남. FA의 online max/lse 유지가 관건 ↗ V07 §4.

3 Online Softmax 누적

4 Causal Load Imbalance ★

• rank r 의 일 = Σ_t≤r |K_t| ∝ r+1
• max rank 가 min rank 대비 P× 일

5 Striped CP balance

원래: rank r = [r·S/P, (r+1)·S/P) 연속 블록
striped: rank r = [r, r+P, r+2P, ...] stride
 → 어느 block 이든 causal workload ≈ uniform
 (앞·뒤 mix)

Context Parallel (NV), "Striped Attention" (Brandon 2023).

6 통신량

7 CP vs Ulysses 비교

8 CP & 다른 축 조합

• CP + TP: attention 내부를 CP, MLP/proj 는 TP (독립 그룹)
• CP + EP: MoE long-context 학습
• CP = DP 처럼 data 축으로 보이지만 attention 에선 sync 필요

SP/CP attention 상세 알고리즘 ↗ V07 §11.

1 전체 조합 수식

2 축별 통신 요약

3 축 배치 원칙 (!) TP≤8 · PP

4 모델 크기별 매트릭스 ★★

예시 구성. DP·PP·TP 곱 = GPU. ZeRO-3 는 DP 내에서 ZeRO-1/2/3 중 선택. 숫자는 대표적 점 — 실제 tuning 필요.

5 7B 선택 근거

• param·grad·opt 모두 한 node (8 GPU) 에 ZeRO-3 로 적재
• TP=1: 작은 모델은 TP 통신 overhead 손해 더 큼
• long-ctx: CP 확장으로 memory 해결

6 70B / 400B 선택 근거

• 70B: TP=8 + DP(ZeRO-3) · node 내에 layer 1개가 다 들어가는 상한
• 400B: TP=8 + PP=8 (node 간 layer 쪼갬) + DP 로 scale-out
• PP 는 bubble (P-1)/M 을 만족할 M 필요 (M ≥ 32 목표)

7 MoE 특수

• EP 축 추가 → AllToAll 비용 (node 내 유지가 최선)
• TP 줄이고 EP 확장 경향 (expert 가 이미 param 분산 역할)
• DeepSeek V3: TP=1, PP 로 파이프 · EP=8 내부

8 Rank 매핑 공식

9 Node affinity

실무 rule of thumb: TP·CP·EP 는 intra-node, DP·PP 는 inter-node. AllToAll 비용을 NVLink 안에 가둔다.

1 Overlap 기본 원리

• SM 소비: NCCL kernel 은 channel × SM 을 얕게 잡음 (~2-4 SM)
• 나머지 SM 은 compute 여유 → overlap 성립

2 Backward AllReduce Bucketing

naive: grad 1개 생성마다 AR → 작은 msg × 많음
bucket: 연속 param grad 를 buffer 에 쌓아 합치기
 └─ 25 MB 정도 버킷 권장
 └─ 첫 bucket ready 시점부터 AR 시작
 └─ 이후 bucket 은 compute 와 overlap

PyTorch DDP bucket_cap_mb=25 기본. bucket 클수록 BW 효율 ↑, latency 시작 ↑.

3 FSDP Prefetch ★

layer_n forward  :  AG(param_n)   ┐
                    compute_n      │
layer_n+1 forward: AG(param_n+1) prefetch (overlap)

layer_n backward : AG(param_n)    │
                   compute_n      │
                   RS(grad_n)     │ overlap with next layer's AG
layer_n-1 backward: AG(param_n-1) prefetch

• PyTorch FSDP backward_prefetch=BACKWARD_PRE
• 통신 stream + compute stream 분리 필수

4 Async Collective API

• NCCL: kernel launch 후 반환 → CPU 는 다음 작업
• PyTorch async_op=True → work handle
• 수동 overlap: handle.wait() 를 compute 이후에 배치

5 PP & overlap

• 1F1B 의 backward 기간 동안 다음 forward 와 P2P 교차
• stream 2개: activation P2P · grad P2P 분리
• NCCL send/recv 를 서로 다른 comm 으로 → 큐 충돌 회피

6 TP & overlap (advanced)

• Megatron "async tensor parallel" — WGMMA 와 interleave  ↗ V04 §7
• Ring-allreduce 는 natural pipelining

7 Overlap 한계

8 CUDA Graph & overlap

• 전체 iteration 을 graph capture → launch overhead 제거
• decode path 에 효과 극대  ↗ V16 §10
• NCCL 2.11+ 부분 지원 (CUDA_GRAPH_MIXING_SUPPORT)

9 SM 사용량 분배

1 α-β 모델 정의 ★

• α ∝ hop count · driver · NIC 지연
• β ∝ 1/link_BW

2 확장: α-β-γ

BF16 allreduce 에선 γ ≪ β, 무시해도 됨. FP32 + 큰 tensor 에선 γ 가 드러남.

3 Typical 값 H100 DGX

대략적 — driver/stack 에 따라 변동. 출처: NCCL-tests nominal 값.

2 AllReduce 이론 시간

5 Switching point ★

실무: N=8, α=5μs, β=1ns/B → S_* ≈ 수 KB 범위.

6 이론 하한 — Bandwidth

7 이론 하한 — Latency

8 Effective BW 식

9 분해와 α 저감 (!) bucket = α↓

1 RDMA 원리

• registered memory region (MR) 의 rkey 로 원격 접근
• completion queue (CQ) 로 비동기 신호

2 InfiniBand vs RoCE

RoCE v2 는 UDP 위에서 돌기 때문에 IP routing 가능 — 대규모 cloud 에 유리.

3 Queue Pair (QP) ★

• RC (Reliable Connected) · UD (Unreliable Datagram) · UC (Unreliable Connected)
• NCCL : RC 주로 사용, DC (Dynamic Connected) 실험

4 SEND/RECV vs RDMA WRITE ★

NCCL intra-collective 는 주로 RDMA WRITE (receiver 가 어디에 둘지 알고 있음). 큰 msg 에 유리.

5 Work Request 흐름

app → ibv_post_send(wr)   # doorbell PCIe MMIO
NIC ─ DMA from reg-mem ─→ wire
wire ─→ remote NIC ─ DMA to reg-mem ─→ remote
complete → CQ entry → app polls / event

6 GPUDirect RDMA ★ (!) NIC ↔ GPU

• 조건: PCIe ACS off · NIC-GPU 같은 root complex 또는 NVLink
• 성능: latency 2-3× 감소, BW 실효 ↑

7 GPUDirect Storage 구분

GDR (NIC ↔ GPU) 과 GDS (NVMe ↔ GPU) 는 별개. NCCL 은 GDR 만 사용. GDS 는 I/O 용.

8 IB SHARP in-net

• ConnectX-6+ 지원 · subnet manager 설정 필요
• NCCL NCCL_COLLNET_ENABLE=1

9 디바이스 선택 환경

• NCCL_IB_HCA=mlx5_0,mlx5_1 명시적 NIC 지정
• NCCL_IB_GID_INDEX RoCE v2 에서 필수
• NCCL_NET_GDR_LEVEL=PXB GDR 레벨 조정
• NCCL_SOCKET_IFNAME bootstrap 용 TCP iface

1 Debug / 진단

첫 실행 시 NCCL_DEBUG=INFO 필수. 어떤 algo/proto 가 쓰이는지 확인.

2 Algorithm 선택

3 Channel 수

• NCCL_MIN_NCHANNELS, NCCL_MAX_NCHANNELS
• default: topology에 따라 자동
• 너무 크면 small msg 초기화 비용 ↑

4 P2P 제어

level: NVL(NVLink) < PXB(same PCIe switch) < PHB(host bridge) < SYS. 허용 수준 명시.

5 IB / 네트워크

6 Buffer 크기

BUFFSIZE ↑ → large msg pipelining ↑, GPU mem 도 ↑.

7 동작 모드

• NCCL_LAUNCH_MODE=PARALLEL — 모든 rank 동시 launch
• NCCL_ASYNC_ERROR_HANDLING=1 — blocking 대신 이벤트로 error
• NCCL_BLOCKING_WAIT=1 — wait 시 spin

8 Tuning 실마리 ★

9 Reproducibility

1 MPI 위치

• 구현체: OpenMPI · MPICH · MVAPICH · Intel MPI
• GPU-aware MPI: CUDA ptr 을 send 에 직접 전달 가능

2 NCCL 위치

• GPU-GPU 만 다룸 (CPU 쪽 MPI 보완 관계)
• collective + P2P 지원 (2.7+)

3 기능 비교 ★

4 Hybrid 패턴 ★ (!) MPI rank + NCCL comm

MPI_Init → MPI_Comm_rank/size
if rank==0: ncclGetUniqueId(&id)
MPI_Bcast(&id, ...)
cudaSetDevice(local_rank)
ncclCommInitRank(&comm, size, id, rank)
// now NCCL for collectives, MPI for orchestration

5 torchrun 대안

• PyTorch: torchrun·torch.distributed.init_process_group("nccl")
• 내부적으로 TCPStore (rendezvous) + NCCL
• MPI 없이 동작 — 많은 학습 스택에서 기본

6 Launcher 비교

7 MPI 선택 이유

• HPC cluster (Slurm + PMIx) — 기존 infra 재사용
• CPU 측 complex workflow (preprocessing · simulation)
• I/O 에 MPI-IO 필요 (large dataset stream)
• MPI RMA (one-sided) 로 fault handling

8 NCCL 단독 선택

• 순수 학습/추론 스택 — overhead ↓
• cloud k8s 환경 (TCPStore rendezvous 편리)
• elastic/fault-tolerant 중시 (torch elastic)

9 상호 겹침 주의

10 GLOO · UCC 맥락

PyTorch gloo backend: CPU collective. ucc: UCX 기반 hybrid. NCCL 이 안 되는 CPU/TPU 대체용.

1 Parallelism 선택 결정 트리 ★

model & batch 한 GPU fit ?
├─ YES → DP (DDP or ZeRO-1)  · done
└─ NO
   │
   param 한 GPU fit ?
   ├─ YES → ZeRO-2 / ZeRO-3 (FSDP)
   └─ NO
      │
      layer 1개 한 GPU fit ?
      ├─ YES → TP + ZeRO-3
      └─ NO
         │
         layer set 하나가 intra-node 들어감 ?
         ├─ YES → TP + PP + ZeRO
         └─ NO → TP + PP + EP (if MoE) + DP

seq > 32K ?                 → + SP (Ulysses)
seq > 128K ?                → + CP (ring attention)
MoE layer 포함 ?            → + EP

2 Collective × topology BW 표 ★★

bytes/GPU = 양방향 합. step 은 α 계수. NCCL tuner 가 threshold 로 자동 선택.

3 5-축 통신 요약

4 α-β 빠른 추정

5 NVLink intra vs IB inter BW

"typical eff" 는 AllReduce busBW 기준 근사. NCCL-tests 결과와 유사.

6 디버깅 10대 체크

1. NCCL_DEBUG=INFO 로그 확인 — algo/proto
2. nvidia-smi topo -m 으로 GPU-NIC affinity
3. rank ↔ GPU 매핑 (CUDA_VISIBLE_DEVICES)
4. IB HCA 명시: NCCL_IB_HCA
5. PXN 활성 / ACS off 확인
6. bucket size (DDP) · prefetch (FSDP) 설정
7. stream 분리 — 통신 ↔ 연산
8. TP ≤ 8 (intra-node), PP bubble ratio ≤ 0.25
9. GPUDirect RDMA 동작 (peermem loaded)
10. NCCL version 고정 · 여러 버전 혼재 금지

7 관련 권

8 5-축 니모닉 (!) DT-PE-C

9 한 줄 요약

NVLink 는 TP·SP 안에 가두고, IB 는 DP·PP 사이로만 보낸다. AllToAll 은 intra-node 우선. ring 은 large-S BW 왕, tree/NVLS 는 small-S latency 왕.