AEGLOSLAB 엔비디아 기술 스택

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

4.6 NVLink·NVSwitch I — 링크·대역폭

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

4.5에서 NVLink를 "도메인" 수준으로 맥락에서 봤다. 이 장은 그 기초를 체계적으로 정리한다 — NVLink가 물리적으로 무엇이고, 대역폭을 어떻게 세며, GPU가 서로의 메모리에 직접 접근하는 peer-to-peer(P2P) 가 어떻게 동작하는지. 그리고 PCIe와 비교해 왜 다GPU에서 NVLink가 결정적인지를 직접 측정·확인한다. 토폴로지(도메인 형태)는 4.7로 잇는다.

배우는 것:

  • NVLink 링크·대역폭을 세는 법 — per-link × 링크 수 × 방향(양방향), 집계 대역폭 계산.
  • NVLink vs PCIe 대역폭 격차와, 그것이 다GPU 통신·집합연산(4.13)에 주는 영향.
  • P2P 접근(GPU가 다른 GPU HBM을 직접 read/write)과 GPUDirect P2P, 통합 메모리.
  • nvidia-smi topo -m으로 토폴로지를 읽고, P2P 대역폭을 실측한다.

선수 지식: 1.9(HBM·대역폭 — NVLink는 GPU 간 대역폭), 2.5(전송 최적화·PCIe H2D — NVLink는 그 GPU 간 버전), 4.1(폼팩터 — SXM은 풀 NVLink, PCIe는 제한), 4.5(NVLink 도메인 맥락), 4.13(집합연산 — NVLink 위에서 동작). 환경은 0.3. NVLink 세대별 per-link·링크 수·총 대역폭(GB/s)은 세대·시점 의존 → 데이터시트로 확인. 이 장은 세는 법과 측정에 집중한다.

4.6.2 NVLink란 — GPU 간 전용 고속 링크

PCIe는 범용 I/O 버스다 — GPU·NIC·스토리지가 다 같이 쓰며, GPU 간 통신엔 대역·지연이 부족하다. NVLink는 GPU↔GPU(및 GPU↔CPU) 전용 고속 인터커넥트로, PCIe보다 훨씬 큰 대역으로 GPU를 직접 잇는다.

PCIe — 범용·공유, 상대적 저대역 NVLink — GPU 전용, 고대역, peer 메모리 전용 고속 링크 GPU PCIe 스위치/ 루트 CPU / 다른 GPU GPU GPU

용어:

  • 링크(link): NVLink 한 가닥. 각 링크는 여러 lane(차동 신호쌍)으로 구성.
  • GPU당 링크 수: 한 GPU가 가진 NVLink 링크 개수(세대마다 증가).
  • 방향: NVLink는 양방향. 대역폭을 말할 때 단방향(per direction)인지 양방향 합(bidirectional)인지 반드시 구분해야 한다(혼동의 주범).
link0 link1 ... N개 링크 GPU GPU

각 링크는 단방향 X GB/s, 양방향 2X GB/s. GPU당 집계 = per-link 대역 × 링크 수 × (양방향이면 ×2).

4.6.3 대역폭 세는 법

NVLink 집계 대역폭 계산의 골격(수치는 세대 의존이라 변수로):

  • $\mathrm{per\_link\_unidirectional} = L$ (GB/s, 한 링크 한 방향 — 세대 의존)
  • $\mathrm{links\_per\_gpu} = N$ (GPU당 NVLink 링크 수 — 세대 의존)
  • GPU당 단방향 집계 $= L \times N$
  • GPU당 양방향 집계 $= L \times N \times 2$ — 보통 마케팅이 말하는 "NVLink 대역폭"

예시 계산(가상 수치, 실제는 확인 필요). per-link 단방향 L=25 GB/s, GPU당 N=18 링크라면:

  • 단방향 집계 $= 25 \times 18 = 450\,\mathrm{GB/s}$
  • 양방향 집계 $= 450 \times 2 = 900\,\mathrm{GB/s}$

(위 25·18·900은 세는 법을 보이기 위한 예시 숫자다. 실제 세대별 값은 데이터시트로 — 외우지 말 것.)

왜 단/양방향 구분이 중요한가. 벤치마크·논문·마케팅이 같은 링크를 단방향/양방향으로 다르게 인용해 2배 차이가 난다. all-reduce(4.13) 같은 집합연산은 송수신이 동시라 양방향이 의미 있지만, 단방향 복사는 단방향 대역이 한도다. 수치를 비교할 땐 항상 어느 쪽인지 확인.

HBM 대역폭과의 관계. NVLink 집계는 보통 그 GPU의 HBM 대역폭(1.9)보다 작다 — GPU가 자기 HBM을 읽는 것보다 옆 GPU HBM을 NVLink로 읽는 것이 느리다. 그래서 데이터 지역성(자기 HBM 우선)이 여전히 중요하고, GPU 간 통신은 최소화·최적화 대상이다(4.14).

4.6.4 NVLink vs PCIe — 왜 다GPU에서 결정적인가

축NVLinkPCIe
용도GPU↔GPU 전용범용 I/O(GPU·NIC·SSD 공유)
대역폭큼(GPU당 집계 수백 GB/s급, 세대 의존)작음(예 Gen5 x16 양방향 ~수십~백 GB/s급, 확인 필요)
peer 메모리직접·고대역가능하나 느림
지연낮음높음
결합NVSwitch로 전결합(4.7)트리·루트 경유

핵심: 다GPU 학습의 매 스텝 통신(텐서 병렬·all-reduce)이 PCIe로 가면 병목이다 — PCIe 대역이 NVLink보다 한참 작고, 다른 트래픽(NIC·스토리지)과 경쟁하며, 호스트(CPU)를 거치기도 한다. NVLink는 이것을 GPU 직통 고대역으로 해결한다.

다GPU all-reduce(4.13)에서 경로에 따라 확장성이 갈린다:

경로통신 양상확장
PCIe만GPU들이 PCIe·호스트 거쳐 교환 → 대역 부족·경쟁 → 통신이 스텝 지배✗
NVLinkGPU 직통 고대역 → 통신이 연산 뒤에 은닉됨✓

이것이 학습 클러스터가 SXM·NVLink·NVSwitch로 가는 핵심 이유다(4.1).

이것이 4.1의 "SXM(풀 NVLink) vs PCIe(제한)" 차이가 다GPU 성능을 가르는 메커니즘이다.

4.6.5 P2P 접근 — GPU가 GPU 메모리를 직접

NVLink의 핵심 기능은 peer-to-peer(P2P) 메모리 접근 — 한 GPU가 다른 GPU의 HBM을 직접 read/write(호스트를 안 거치고). 이것이 GPUDirect P2P다.

P2P 없음 (PCIe 2번, 느림) P2P 있음 (직통, 호스트 우회) 복사 복사 NVLink GPU0 HBM 호스트 RAM GPU1 HBM GPU0 HBM GPU1 HBM

CUDA에서 P2P 활성화·복사(개략):

// 1) 두 GPU가 P2P 가능한지 확인 후 활성화
int canAccess = 0;
cudaDeviceCanAccessPeer(&canAccess, 0, 1);   // GPU0이 GPU1에 접근 가능?
if (canAccess) {
    cudaSetDevice(0); cudaDeviceEnablePeerAccess(1, 0);  // 0→1
    cudaSetDevice(1); cudaDeviceEnablePeerAccess(0, 0);  // 1→0 (양방향 따로)
}
// 2) P2P 직접 복사 (호스트 우회, NVLink 경유면 고대역)
cudaMemcpyPeer(dst_on_gpu1, 1, src_on_gpu0, 0, nbytes);
// 3) 커널이 peer 포인터를 직접 역참조도 가능 (peer access 켜진 상태)
  • cudaDeviceCanAccessPeer — 두 GPU가 P2P 가능한 연결(NVLink·또는 PCIe P2P)인지.
  • cudaMemcpyPeer — GPU↔GPU 직접 복사. NVLink로 연결됐으면 고대역, PCIe면 느림.
  • peer access가 NVLink를 경유하는지 PCIe를 경유하는지가 성능을 가른다 — 토폴로지(4.7) 확인 필수.

NCCL(4.13)은 이 P2P를 내부에서 활용해 집합연산을 구현한다 — 우리가 직접 cudaMemcpyPeer를 쓸 일은 드물고, 보통 NCCL·프레임워크가 NVLink P2P를 자동 사용.

4.6.6 워크스루 — 토폴로지 확인·P2P 측정

1) 토폴로지 보기 — nvidia-smi topo -m. GPU 간 연결 종류를 행렬로 출력:

nvidia-smi topo -m
# 출력(개략 예):
#        GPU0  GPU1  GPU2  GPU3
# GPU0    X    NV#   NV#   SYS
# GPU1   NV#    X    SYS   NV#
# ...
# 범례: NV# = NVLink(숫자=링크 수), PIX/PXB/PHB = PCIe 경로, SYS = 시스템(소켓 간) 경로

읽는 법:

  • NV# (예 NV6) — 두 GPU가 NVLink로 연결, 숫자가 링크 수(대역 가늠). 이것이 보이면 고대역 P2P.
  • SYS·PHB·PXB·PIX — PCIe 경로(호스트/소켓 거침), 대역 낮음. 이 쌍의 P2P는 느리다.
  • 다GPU 작업 전 이 행렬로 어느 GPU 쌍이 NVLink로 빠르고 어느 쌍이 PCIe로 느린지 파악 → 통신 집약 작업을 NVLink 쌍에 배치(4.7·4.14).

2) P2P 대역폭 실측. CUDA 샘플의 p2pBandwidthLatencyTest(또는 직접 cudaMemcpyPeer 반복 측정):

# CUDA Samples 빌드 후
./p2pBandwidthLatencyTest
# 출력(개략): GPU 쌍별 P2P 대역폭(GB/s)·지연(us) 행렬
#   - NVLink 쌍: 높은 GB/s (수백 GB/s급, 세대 의존)
#   - PCIe 쌍:   낮은 GB/s
#   - P2P 켰을 때 vs 껐을 때(호스트 경유) 비교도

관찰: NVLink로 연결된 쌍은 P2P 대역이 PCIe 쌍보다 훨씬 높다. P2P를 끄면(호스트 경유) 대역이 더 떨어진다. 이 실측이 "왜 NVLink가 중요한가"를 숫자로 보여준다 — 토폴로지(4.7)와 병렬화 매핑(4.14)의 근거.

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

  • [ ] NVLink 집계 대역폭을 per-link × 링크 수 × 방향으로 세고, 단방향/양방향을 구분한다.
  • [ ] NVLink와 PCIe의 대역 격차가 왜 다GPU 통신을 가르는지(매 스텝 통신·경쟁·호스트 우회) 설명한다.
  • [ ] P2P 접근(GPUDirect P2P)이 호스트를 우회해 GPU↔GPU 직통함을 알고, NVLink 경유 vs PCIe 경유 차이를 안다.
  • [ ] nvidia-smi topo -m을 읽어 NVLink 쌍과 PCIe 쌍을 구분하고, P2P 대역을 실측할 줄 안다.

증상→의심: 다GPU 학습 통신 병목·확장 안 됨 → GPU 쌍이 NVLink 아닌 PCIe 연결(topo -m 확인). P2P 복사가 느림 → NVLink 아닌 경로(SYS/PHB) 또는 P2P 미활성. 대역폭 수치가 문서와 2배 차이 → 단/양방향 혼동. peer access 활성화 실패 → 그 쌍이 P2P 불가 연결(토폴로지·드라이버). NCCL이 느림 → NVLink 미사용(토폴로지·NCCL 설정, 4.13).

4.6.8 트레이드오프·맥락

연결대역용도
자기 HBM(1.9)최대항상 우선 — 지역성
NVLink(P2P)큼노드 내 GPU 간(스케일업)
PCIe(P2P)작음NVLink 없는 GPU 쌍·범용
InfiniBand(4.11)노드 간노드를 넘는 스케일아웃

견해: 대역폭 위계(자기 HBM > NVLink > PCIe > IB)를 항상 기억하라. 최적화의 방향은 통신을 위계 위쪽(빠른 쪽)에 가두는 것 — 자기 HBM에서 끝낼 수 있으면 최고, 안 되면 NVLink 도메인 안에서(4.5), 그래도 안 되면 IB로. 직접 cudaMemcpyPeer를 쓰는 일은 드물고, NCCL·프레임워크가 토폴로지를 보고 NVLink를 자동 활용하게 두되, 토폴로지가 기대대로인지(topo -m) 확인하는 것이 우리 몫.

비NVIDIA 정직: GPU 전용 고대역 인터커넥트·peer 메모리는 NVLink만의 개념이 아니다 — 다른 가속기도 유사 링크를, 업계는 이더넷 기반 스케일업/아웃 표준(Ultra Ethernet·UALink 등 흐름)을 추진한다. NVLink의 강점은 성숙도·대역·NVSwitch 전결합·NCCL 통합. 사양은 세대마다 급변하니 세는 법(이 장)을 익혀 그때그때 데이터시트로 계산·비교하라.

4.6.9 흔한 함정

  • 단/양방향 혼동. "NVLink 900 GB/s"가 양방향 합인지 단방향인지 모르고 비교하면 2배 틀린다. 항상 명시·확인.
  • PCIe 연결을 NVLink로 가정. topo -m에서 SYS/PHB인 쌍을 NVLink로 착각하면 통신 병목의 원인을 놓친다. 토폴로지를 실제로 확인.
  • P2P 미활성 방치. peer access를 안 켜면 GPU 간 복사가 호스트를 거쳐 느리다. NCCL은 자동이지만 직접 코드는 cudaDeviceEnablePeerAccess 확인.
  • NVLink 대역을 HBM 대역과 동일시. NVLink는 HBM(자기 메모리)보다 느리다 — 옆 GPU 데이터를 NVLink로 읽는 것은 자기 HBM보다 비용이 크다. 지역성 우선.
  • 토폴로지 무시한 GPU 배치. 통신 집약 작업을 PCIe로만 연결된 GPU 쌍에 배치하면 NVLink가 있어도 못 쓴다. 작업을 NVLink 쌍/도메인에 매핑(4.7·4.14).
  • 클라우드에서 토폴로지 미확인. 인스턴스마다 GPU 간 연결이 다르다(풀 NVLink vs PCIe). topo -m으로 실제 확인 후 병렬화 결정.
  • 마케팅 수치 단정. per-link·링크 수·집계는 세대·시점 의존. 세는 법만 외우고 값은 데이터시트로.

4.6.10 연습문제

연습 1 (대역폭 계산). per-link 단방향 L=50 GB/s, GPU당 링크 수 N=12인 가상 세대. (a) 단방향 집계, (b) 양방향 집계를 구하고, (c) 문서가 "이 GPU NVLink 1200 GB/s"라 했다면 단/양방향 중 무엇일지 추론하라.

정답

(a) 단방향 집계 $= 50 \times12 = 600 \,\mathrm{GB/s}$. (b) 양방향 집계 $= 600 \times2 = 1200 \,\mathrm{GB/s}$. (c) 1200은 양방향 집계(=600×2)와 일치하므로 양방향 합으로 인용한 것. 단방향 작업(예 한 방향 복사)의 한도는 600 GB/s임에 주의. 마케팅은 보통 양방향 합을 쓴다 — 비교 시 통일.

연습 2 (토폴로지 해석). nvidia-smi topo -m에서 GPU0–GPU1이 NV6, GPU0–GPU2가 SYS로 나왔다. (a) 두 쌍의 P2P 성능 차이, (b) 텐서 병렬(통신 집약) 2-GPU 작업을 어느 쌍에 배치해야 하나, (c) 왜.

정답

(a) GPU0–GPU1(NV6)은 NVLink 6링크로 고대역 P2P, GPU0–GPU2(SYS)는 시스템(PCIe/소켓 간) 경로로 저대역·고지연 P2P. (b) GPU0–GPU1에 배치한다. (c) 텐서 병렬은 매 레이어 막대한 GPU 간 통신(all-reduce)이라 NVLink 고대역이 필수 — SYS 쌍에 배치하면 PCIe·호스트 경유로 통신이 병목 되어 확장이 안 된다. 토폴로지를 보고 통신 집약 작업을 NVLink 쌍에 배치(4.14).

연습 3 (P2P 경로). GPU0이 GPU1의 배열을 더하는 커널을 짠다. (a) P2P(NVLink)로 직접 읽을 때와 (b) 호스트를 거쳐 복사 후 읽을 때의 데이터 경로·대역 차이를 그리고, P2P 활성화 코드 한 줄을 쓰라.

정답

(a) P2P/NVLink: GPU1 HBM ──NVLink──▶ GPU0(커널이 peer 포인터 직접 역참조, 호스트 우회, 고대역). (b) 호스트 경유: GPU1 HBM →(PCIe) 호스트 RAM →(PCIe) GPU0 HBM(PCIe 2회, 저대역·고지연, CPU 메모리 왕복). 차이: P2P는 PCIe·호스트를 안 거쳐 훨씬 빠름(NVLink 연결 시). 활성화: cudaSetDevice(0); cudaDeviceEnablePeerAccess(1, 0);(GPU0이 GPU1 접근 허용; 양방향이면 반대도). 사전에 cudaDeviceCanAccessPeer로 가능 여부 확인.

연습 4 (위계 적용). "통신을 대역폭 위계 위쪽에 가둔다"는 원칙을, 한 모델을 (a) 자기 HBM, (b) NVLink 도메인, (c) 노드 간 IB로 단계적으로 펼치는 의사결정으로 설명하라.

정답

대역 위계: 자기 HBM > NVLink > IB. 원칙은 통신을 가능한 빠른 계층에 가두는 것. (a) 모델·상태가 한 GPU에 들어가면 자기 HBM에서 끝 — GPU 간 통신 0(최선). (b) 한 GPU를 넘으면 NVLink 도메인 안(4.5)에 펼쳐 통신 집약 병렬화(텐서·전문가)를 NVLink 고대역으로. (c) 도메인도 넘으면 노드 간 IB(4.11)로 — 단 IB는 느리니 통신이 드문 병렬화(데이터·파이프라인)를 이 경계에 매핑(4.14). 즉 통신량이 큰 병렬화일수록 위계 위쪽에, 모델이 커질수록 아래로 단계적 확장.

4.6.11 요약 · 다음 장

  • NVLink는 GPU↔GPU(↔CPU) 전용 고속 인터커넥트. 집계 대역은 per-link × 링크 수 × 방향으로 세며, 단방향/양방향을 반드시 구분(2배 혼동 주의). 수치는 세대 의존이니 세는 법을 익히고 데이터시트로.
  • NVLink는 PCIe보다 대역이 한참 커, 다GPU의 매 스텝 통신을 PCIe 병목에서 구한다 — SXM·NVLink가 학습 클러스터의 토대(4.1).
  • P2P(GPUDirect P2P) 로 GPU가 다른 GPU HBM을 호스트 우회해 직접 read/write. NVLink 경유면 고대역, PCIe 경유면 느림 — 토폴로지가 성능을 가른다.
  • 실무: nvidia-smi topo -m으로 NVLink/PCIe 쌍을 확인하고 P2P 대역을 실측해, 통신 집약 작업을 NVLink 쌍/도메인에 매핑. 대역 위계(HBM>NVLink>PCIe>IB)에 통신을 가둔다.

다음 장: 4.7 NVLink·NVSwitch II — 도메인 토폴로지. 링크·대역(이 장)을 여러 GPU의 연결 형태로 확장한다. NVSwitch가 GPU를 전결합하는 방식, 노드 내 8-GPU 토폴로지, 그리고 NVLink 도메인이 노드를 넘어 커질 때(NVL72) 토폴로지가 통신·병렬화에 주는 영향을 본다.