AEGLOSLAB 엔비디아 기술 스택

제4부 데이터센터와 스케일

4.13 NCCL I — 집합연산 원리·알고리즘

4.13.1 이 장에서 배우는 것 / 선수 지식

4.6~4.12에서 GPU를 잇는 하드웨어(NVLink·NVSwitch·IB·GPUDirect)를 쌓았다. 그 위에서 실제로 분산 통신을 수행하는 라이브러리가 NCCL(NVIDIA Collective Communications Library, "니클") 이다. 분산 학습의 거의 모든 GPU 간 통신 — 데이터 병렬의 그래디언트 동기, 텐서·전문가 병렬의 교환 — 이 NCCL의 집합연산(collective) 으로 내려간다. 이 장은 집합연산이 무엇이고, ring·tree 알고리즘이 토폴로지·대역폭에 어떻게 매핑되는지, 왜 ring all-reduce가 대역폭 최적인지를 실습으로 익힌다.

배우는 것:

  • 핵심 집합연산 — all-reduce·reduce-scatter·all-gather·broadcast·all-to-all — 의 정의와 쓰임.
  • all-reduce = reduce-scatter + all-gather 분해와, ring all-reduce가 왜 대역폭 최적(GPU 수에 무관한 ~2S)인지 수계산.
  • ring(대역폭 최적, 지연 ~N) vs tree(지연 ~log N) 트레이드오프와 NCCL의 선택.
  • NCCL이 토폴로지(NVLink·IB·GPUDirect·SHARP)에 매핑되는 방식과 ncclAllReduce·nccl-tests 사용.

선수 지식: 4.6~4.7(NVLink·토폴로지·bisection), 4.11~4.12(네트워크·GPUDirect RDMA·SHARP — NCCL이 그 위에서 동작), 4.5(NVLink 도메인), 4.14(병렬화 매핑 — 다음 장, NCCL의 응용), 6.3(NeMo 병렬화). 환경은 0.3. 코드는 NCCL·CUDA 12.x 가정, 정확한 API는 버전 확인. 대역·성능은 토폴로지·시점 의존.

4.13.2 집합연산이란 — 분산 학습의 통신 단위

집합연산(collective) 은 여러 GPU가 함께 수행하는 통신 패턴이다. 점대점(한 GPU→다른 GPU)이 아니라 그룹 전체가 참여한다. 분산 학습의 통신은 거의 다 이 패턴으로 환원된다.

핵심 집합연산(N개 GPU, 각자 데이터 보유):

집합연산동작
all-reduce모든 GPU의 데이터를 reduce(합산 등)해 → 결과를 모든 GPU가 가짐 (데이터 병렬 그래디언트 동기의 핵심)
reduce-scatterreduce한 결과를 조각내 → 각 GPU가 자기 조각만 가짐
all-gather각 GPU의 조각을 모아 → 모든 GPU가 전체를 가짐
broadcast한 GPU의 데이터를 → 모든 GPU에 복사
reduce모든 GPU의 데이터를 reduce해 → 한 GPU만 결과
all-to-all각 GPU가 모든 GPU에 서로 다른 조각을 전송 (MoE 전문가 라우팅, 4.5)
전 (각자 값 보유) 후 (모두 같은 합) all-reduce 합산 GPU0: a GPU1: b GPU2: c GPU3: d GPU0: a+b+c+d GPU1: a+b+c+d GPU2: a+b+c+d GPU3: a+b+c+d

all-reduce(합산, GPU 4개, 각자 값 보유): 합산 전 각 GPU가 서로 다른 값을 갖고, all-reduce 후 모든 GPU가 같은 합 a+b+c+d를 가진다.

왜 all-reduce가 중심인가. 데이터 병렬 학습(6.3)은 각 GPU가 다른 데이터로 그래디언트를 계산한 뒤, 모든 GPU의 그래디언트를 평균(합산/N) 해 같은 갱신을 적용한다 — 정확히 all-reduce다. 스텝마다 일어나므로 all-reduce 효율이 분산 학습 속도를 좌우한다.

4.13.3 all-reduce = reduce-scatter + all-gather

순진한 all-reduce(모두가 모두에게 보내 각자 합산)는 통신량이 폭발한다($O(N^2)$ 또는 한 GPU에 집중). NCCL은 두 단계로 분해해 효율적으로 한다:

all-reduce reduce-scatter (각자 다른 조각을 합산해 나눠 가짐) all-gather (그 합산 조각들을 모두가 모음)
데이터를 N조각으로 나눔 (S/N each) GPU 4개, 조각 0·1·2·3 ① reduce-scatter: GPU_i가 '조각 i의 전체 합'을 책임지고 모음 → 끝나면 GPU0=합(조각0)·GPU1=합(조각1)·… (각자 다른 합산 조각) ② all-gather: 각자의 합산 조각을 서로 실행해 모두가 전체 합을 가짐 → 끝나면 모든 GPU가 합(0)·합(1)·합(2)·합(3) 전체 = all-reduce 결과

이 분해가 ring 알고리즘의 토대다 — 각 단계를 고리로 실행하면 통신이 모든 링크에 고르게 분산돼 대역폭 최적이 된다.

4.13.4 ring all-reduce — 왜 대역폭 최적인가

ring 알고리즘: GPU들을 고리(ring) 로 잇고, 데이터를 조각 단위로 한 방향으로 흘린다. 각 GPU는 동시에 한 조각을 보내고 한 조각을 받는다(이웃에게서) — 모든 링크가 동시에 일한다.

GPU0 GPU1 GPU2 GPU3

고리(ring): 각 GPU가 매 스텝 왼쪽 이웃에서 받고 + 오른쪽 이웃에 보낸다(동시, 모든 링크 활성).

대역폭 수계산(핵심). N개 GPU, 데이터 크기 S. 데이터를 N조각(S/N)으로 나눈다.

  • reduce-scatter: N-1 스텝. 각 스텝에 GPU가 S/N 조각을 보냄. → GPU당 송신 $= (N-1)\cdot S/N$.
  • all-gather: N-1 스텝. 각 S/N. → GPU당 송신 $= (N-1)\cdot S/N$.
  • 합계 GPU당 송신 $= 2\cdot(N-1)\cdot S/N = 2\cdot S\cdot(N-1)/N$.
N이 커지면 (N-1)/N → 1 따라서 GPU당 송신량 → 2·S (N에 무관!) ★ ring all-reduce의 통신량은 GPU 수와 무관하게 ~2S → 대역폭 최적

각 GPU는 자기 데이터를 대략 한 번 내보내고 한 번 받아오는 셈이다.

왜 이것이 놀라운가: GPU를 4개→100개로 늘려도 GPU당 통신량은 ~2S로 거의 일정하다. 통신이 모든 링크에 고르게 퍼져, 링크 대역폭을 최대한 쓴다. 그래서 ring은 대역폭 최적(bandwidth-optimal).

대가 — 지연. 단계 수가 2(N-1)로 GPU 수에 비례. N이 크면 단계가 많아져 지연(latency)이 늘고, 작은 메시지에선 이 단계별 오버헤드가 지배한다. → 큰 메시지엔 ring(대역폭), 작은 메시지·큰 N엔 tree(지연).

all-reduce 시간 ≈ $2\cdot S\cdot(N-1)/N / B$ (대역폭 항) + $2(N-1)\cdot\alpha$ (지연 항):

항식변수지배 조건
대역폭 항2·S·(N-1)/N / BB = 링크 대역큰 S → ring 유리
지연 항2(N-1)·αα = 스텝당 지연작은 S·큰 N → tree 유리

4.13.5 tree·hierarchical·SHARP

tree 알고리즘. GPU를 트리로 묶어 계층적으로 reduce(잎→뿌리)하고 broadcast(뿌리→잎). 단계 수가 O(log N)이라 지연이 GPU 수에 로그로만 증가 — 큰 클러스터·작은 메시지에 유리. 대신 대역폭 효율은 ring보다 낮을 수 있다.

알고리즘단계 수특성
ring~N대역폭 최적, 큰 메시지
tree~log N지연 최적, 작은 메시지·큰 N
NCCL—메시지 크기·GPU 수·토폴로지를 보고 ring/tree(및 변형)를 자동 선택

hierarchical(계층) 매핑. 실제 클러스터는 NVLink(노드 내, 빠름) + IB(노드 간, 느림)의 2계층(4.5~4.11). NCCL은 이를 계층적으로 처리한다:

① 노드 내: NVLink로 GPU들 reduce-scatter (고대역) ② 노드 간: IB로 조각별 노드 간 그룹이 all-reduce (느린 링크 트래픽 최소화) ③ 노드 내: NVLink로 all-gather 배포

느린 노드 간(IB) 트래픽을 줄이고 빠른 NVLink를 최대 활용한다(토폴로지 인식).

SHARP 오프로드(4.11). IB 스위치가 in-network reduction을 지원하면, NCCL이 all-reduce의 합산을 스위치에 오프로드 — 노드 간 트래픽·홉을 더 줄인다. 대규모 데이터 병렬에서 큰 이득.

핵심: NCCL은 토폴로지를 자동 탐지(NVLink·IB·친화도)해 ring/tree·계층·SHARP를 조합한다. 우리는 보통 ncclAllReduce 한 번 부르고, NCCL이 최적 경로를 구성한다 — 단 토폴로지·환경변수가 맞아야(4.13.8).

4.13.6 워크스루 — NCCL 호출·측정

1) NCCL all-reduce 골격(개략).

#include <nccl.h>
// 다GPU(또는 다프로세스) 각자: communicator 초기화
ncclComm_t comm;
ncclCommInitRank(&comm, nRanks, ncclId, myRank);   // nRanks개 중 myRank
// ... 각 GPU에 grad 버퍼(크기 count, float) ...
ncclAllReduce(sendbuf, recvbuf, count, ncclFloat, ncclSum, comm, stream);
//  → 모든 rank의 sendbuf를 합산해 recvbuf에 (모두 같은 결과). 스트림에 비동기 enqueue
cudaStreamSynchronize(stream);
// 실제로는 PyTorch DDP·NeMo(6.3)가 내부에서 이것을 호출 — 직접 작성할 일은 드묾

PyTorch에선 torch.distributed(백엔드 nccl)가 이를 감싼다 — DDP가 backward 중 그래디언트 all-reduce를 자동 호출하고 연산과 오버랩(2.10).

2) 대역폭 실측 — nccl-tests.

# nccl-tests 빌드 후 (다GPU/다노드)
./build/all_reduce_perf -b 8 -e 1G -f 2 -g 8   # 8B~1GB, 8 GPU
# 출력(개략): 메시지 크기별 algbw(알고리즘 대역)·busbw(버스 대역)·시간
#   작은 크기: 지연 지배(busbw 낮음) — tree가 유리한 영역
#   큰 크기:   대역폭 지배(busbw가 링크 대역에 근접) — ring 영역
#   busbw가 NVLink/IB 대역에 근접하면 건강, 한참 낮으면 토폴로지·경로 문제

3) 진단(4.12 연계).

NCCL_DEBUG=INFO ./all_reduce_perf ...   # 어떤 알고리즘(Ring/Tree)·전송(NVLink/IB/GDRDMA)인지
NCCL_ALGO=Ring   # 또는 Tree — 알고리즘 강제(실험·디버깅)
NCCL_DEBUG_SUBSYS=NET,GRAPH   # 경로·토폴로지 그래프 상세

관찰: busbw가 토폴로지 대역(NVLink/IB)에 근접하면 NCCL이 잘 매핑된 것. 한참 낮으면 GPUDirect 미사용(4.12)·친화도·비균일 토폴로지(4.7)를 의심.

4.13.7 중간 점검 (체크포인트)

  • [ ] all-reduce·reduce-scatter·all-gather·all-to-all을 정의하고, all-reduce가 데이터 병렬 그래디언트 동기임을 안다.
  • [ ] all-reduce = reduce-scatter + all-gather 분해를 그린다.
  • [ ] ring all-reduce의 GPU당 통신량 2S(N-1)/N → ~2S를 유도하고, 대역폭 최적·지연 ~N을 설명한다.
  • [ ] ring(큰 메시지) vs tree(작은 메시지·큰 N)와 NCCL의 계층·SHARP 매핑을 안다.

증상→의심: 다GPU 학습이 확장 안 됨 → all-reduce 통신 바운드(nccl-tests busbw·4.14 매핑). busbw가 링크 대역보다 한참 낮음 → GPUDirect 미사용·친화도(4.12)·비균일 토폴로지(4.7). 작은 모델인데 통신 느림 → 지연 지배(tree·메시지 합치기). 노드 간만 느림 → IB·SHARP 경로 점검(4.11).

4.13.8 트레이드오프·맥락

알고리즘강점약점적합
ring대역폭 최적(~2S)지연 ~N큰 메시지(큰 모델 그래디언트)
tree지연 ~log N대역폭 효율↓작은 메시지·큰 클러스터
계층(NVLink+IB)느린 링크 최소화—다노드(항상)
SHARP 오프로드노드 간 트래픽↓스위치 의존대규모 data-parallel

견해: 알고리즘은 NCCL이 자동 선택하니 직접 고를 일은 드물다 — 우리 몫은 NCCL이 좋은 경로를 쓰도록 환경을 갖추는 것: GPUDirect RDMA·친화도(4.12), 균일 토폴로지(4.7), SHARP 가용성(4.11), 그리고 통신을 연산과 오버랩(2.10, DDP가 자동). nccl-tests로 busbw가 토폴로지 대역에 근접하는지 확인하고, 안 되면 환경(드라이버·peermem·NCCL_*)을 점검. 알고리즘 강제(NCCL_ALGO)는 디버깅·특수 상황에만.

비NVIDIA 정직: ring all-reduce·tree reduction은 NCCL 고유가 아니라 분산 컴퓨팅의 고전 알고리즘(MPI가 수십 년 사용). ring all-reduce의 대역폭 최적성도 잘 알려진 결과다. NCCL의 가치는 이를 NVIDIA 토폴로지(NVLink·IB·GPUDirect·SHARP)에 자동 매핑·최적화한 것. 같은 자리를 다른 라이브러리(oneCCL·RCCL·MPI)가 다른 가속기에서 채운다 — 원리(이 장)는 공통이다.

4.13.9 흔한 함정

  • 순진한 all-reduce 직접 구현. "모두가 한 GPU에 보내 합산 후 broadcast"는 그 GPU·링크에 병목 집중. ring/tree(NCCL)가 통신을 분산해 훨씬 빠르다 — 직접 작성하지 말고 NCCL.
  • busbw를 algbw와 혼동. nccl-tests의 algbw(데이터/시간)와 busbw(링크가 실제로 옮긴 양 기준)는 다르다. all-reduce는 busbw ≈ 2·algbw·(N-1)/N. 토폴로지 대역과 비교할 땐 busbw.
  • GPUDirect·친화도 미설정으로 느린 NCCL. busbw가 낮으면 호스트 bounce·먼 NIC(4.12). NCCL_DEBUG로 GDRDMA·NVLink 사용 확인, peermem·친화도 점검.
  • 작은 메시지를 ring으로. 작은 텐서를 자주 all-reduce하면 지연(스텝 ~N)이 지배. 그래디언트 버킷팅(여러 텐서 합쳐 한 번에, DDP 자동)·tree로 완화.
  • 통신을 연산과 오버랩 안 함. all-reduce를 backward와 겹치지 않으면(2.10) 통신이 그대로 스텝에 더해진다. DDP·프레임워크의 오버랩을 켜고 측정에서 확인.
  • 비균일 토폴로지 무시. NCCL이 ring을 구성할 때 PCIe 쌍(4.7)이 끼면 느려진다. 균일 전결합(SXM)이 아니면 토폴로지·NCCL_* 조정.
  • measurement에 워밍업 포함. 첫 NCCL 호출은 communicator·버퍼 셋업으로 느리다. 워밍업 후 측정(2.13).

4.13.10 연습문제

연습 1 (집합연산 매핑). 다음 분산 학습 통신을 어떤 집합연산으로 구현하는지: (a) 데이터 병렬 그래디언트 평균, (b) 모델 파라미터를 모든 GPU에 초기 배포, (c) MoE에서 토큰을 전문가 GPU로 라우팅, (d) ZeRO에서 샤딩된 파라미터를 forward 전 모음.

정답

(a) all-reduce(합산 후 /N로 평균) — 모든 GPU가 같은 평균 그래디언트를 가져 같은 갱신. (b) broadcast — 한 GPU(또는 rank0)의 초기 파라미터를 모두에게 복사. (c) all-to-all — 각 GPU가 자기 토큰을 해당 전문가가 있는 GPU로, 서로 다른 조각을 교환(4.5). (d) all-gather — 각 GPU가 가진 파라미터 샤드를 모아 전체 파라미터를 복원(ZeRO/FSDP, 6.3). (역으로 그래디언트는 reduce-scatter로 샤딩.)

연습 2 (ring 대역폭). GPU 8개, 각 그래디언트 S=400MB, 링크 단방향 대역 B=100 GB/s(가정). ring all-reduce의 (a) GPU당 총 송신량, (b) 대역폭 항 시간(지연 무시), (c) GPU를 16개로 늘리면 GPU당 송신량은?

정답

(a) $2\cdot S\cdot(N-1)/N = 2\cdot400\cdot7/8 = 700 \,\mathrm{MB}$. (b) 시간 ≈ 송신량/B = 0.7GB / 100GB/s = 7 ms(대역폭 항). (c) N=16: $2\cdot400\cdot15/16 = 750 \,\mathrm{MB}$ — 거의 안 늘어난다(700→750MB). N→∞면 $2\cdot400=800\mathrm{MB}$로 수렴. 이것이 ring의 대역폭 최적성: GPU를 2배로 늘려도 GPU당 통신량이 ~2S로 거의 일정 → 통신량이 GPU 수에 무관(단 단계 수=지연은 N에 비례).

연습 3 (ring vs tree). 같은 클러스터에서 (a) 큰 LLM의 700MB 그래디언트 all-reduce, (b) 작은 32KB 텐서를 1024 GPU에서 all-reduce. 각각 ring과 tree 중 무엇이 유리한지와 이유.

정답

(a) 700MB·큰 메시지 → ring. 시간이 대역폭 항(데이터/대역)에 지배되고 ring은 통신량 ~2S로 대역폭 최적. 지연 항(스텝 수)은 큰 메시지에선 상대적으로 작다. (b) 32KB·1024 GPU → tree. 메시지가 작아 데이터 전송 시간보다 단계별 지연이 지배. ring은 단계 ~N=1024로 지연이 폭발하지만, tree는 ~log₂1024=10단계라 지연이 훨씬 작다. 작은 메시지·큰 N에선 tree의 로그 지연이 결정적. NCCL이 메시지 크기·N을 보고 자동 선택한다.

연습 4 (계층 매핑). 4노드 × 8GPU = 32 GPU 데이터 병렬에서, NCCL이 NVLink(노드 내)와 IB(노드 간)를 어떻게 계층적으로 써 노드 간 트래픽을 최소화하는지 단계로 설명하라.

정답

① 노드 내 reduce-scatter(NVLink): 각 노드 8 GPU가 NVLink 고대역으로 그래디언트를 reduce-scatter → 각 GPU가 부분합 조각을 가짐(느린 IB 안 씀). ② 노드 간 all-reduce(IB): 같은 조각을 맡은 노드 간 GPU들끼리만 IB로 all-reduce → 각 노드의 그 GPU가 조각의 전역 합을 가짐. 노드 간 트래픽이 전체가 아니라 조각·대표 그룹으로 최소화(SHARP면 스위치가 합산 오프로드, 4.11). ③ 노드 내 all-gather(NVLink): 조각별 전역 합을 노드 안에서 NVLink로 배포 → 모든 GPU가 최종 all-reduce 결과. 핵심: 느린 IB는 꼭 필요한 노드 간 트래픽만, 빠른 NVLink로 노드 내를 처리 → 토폴로지 인식으로 노드 간 병목 완화.

4.13.11 요약 · 다음 장

  • NCCL은 NVIDIA 토폴로지 위에서 집합연산(all-reduce·reduce-scatter·all-gather·broadcast·all-to-all)을 구현한다. all-reduce가 데이터 병렬 그래디언트 동기의 핵심.
  • all-reduce = reduce-scatter + all-gather. ring all-reduce는 GPU당 통신량이 ~2S로 GPU 수에 무관(대역폭 최적) — 단 단계 수가 ~N이라 지연은 N에 비례.
  • ring(큰 메시지) vs tree(작은 메시지·큰 N) 를 NCCL이 자동 선택하고, 계층(NVLink+IB)·SHARP 오프로드로 느린 노드 간 트래픽을 최소화한다.
  • 직접 알고리즘을 고르기보다 NCCL이 좋은 경로를 쓰도록(GPUDirect·친화도·균일 토폴로지·오버랩) 환경을 갖추고, nccl-tests busbw로 검증.

다음 장: 4.14 NCCL II — 병렬화 매핑·튜닝(실습). 집합연산(이 장)을 실제 분산 학습 병렬화 전략(데이터·텐서·파이프라인·전문가 병렬)에 매핑한다. 어떤 병렬화가 어떤 집합연산·통신량을 만들고, 그것을 토폴로지(NVLink 도메인·IB)에 어떻게 배치·튜닝하는지를 실습으로 익힌다.