AWS 기술 블로그

분산 트레이닝 관점에서의 AWS 인터커넥트 기술 소개 – AWS 환경에서 NCCL을 이용한 GPU 간 통신

지난 블로그에서는 ENI(Elastic Network Interface)의 구조와 역할, 그리고 p5.48xlarge와 p6-b300 인스턴스에서 EFA(Elastic Fabric Adapter)를 실제로 어떻게 구성하는지 살펴보았습니다. 이번 블로그에서는 이렇게 구성된 EFA 네트워크 위에서 실제 GPU 간 통신이 어떻게 이루어지는지, 그 핵심에 있는 NCCL(NVIDIA Collective Communications Library) 통신에 대해 소개하고자 합니다.

EFA가 고속도로를 깔아주는 인프라라면, NCCL은 그 위에서 수백 대의 GPU가 효율적으로 데이터를 주고받을 수 있도록 교통을 조율하는 관제 시스템에 해당합니다. 아무리 빠른 네트워크가 있어도 통신 패턴이 비효율적이면 성능을 제대로 끌어낼 수 없습니다. 이 블로그를 통해 NCCL이 무엇인지, 어떻게 동작하는지, 그리고 AWS EFA 환경에서 어떻게 최적화되는지를 이해할 수 있기를 기대합니다.

시리즈 블로그 보기

분산 학습과 GPU 간 통신의 필요성

대규모 딥러닝 모델을 학습할 때는 단일 GPU로는 충분하지 않습니다. 수백억 개의 파라미터를 가진 대규모 언어 모델이나 복잡한 이미지 생성 모델을 학습하려면, 수십 개에서 수천 개 이상의 GPU를 동시에 활용하는 분산 학습이 필수적입니다. 분산 학습에서는 각 GPU가 데이터의 일부를 처리하거나 모델의 일부를 담당하며, 학습 과정에서 GPU들이 지속적으로 정보를 교환해야 합니다. 예를 들어, 데이터 병렬 학습에서는 각 GPU가 서로 다른 배치(batch)의 데이터로 그래디언트(gradient, 각 파라미터를 어느 방향으로 얼마나 조정해야 하는지 알려주는 수정 지침서)를 계산한 후, 이를 모든 GPU에 동기화하여 일관된 모델 업데이트를 수행합니다. 이러한 동기화 과정이 매 학습 스텝마다 발생하므로, GPU 간 통신의 효율성이 전체 학습 성능을 좌우하게 됩니다. 수백 개의 GPU를 사용하는 대규모 클러스터에서 통신 병목이 발생하면 전체 학습 시간이 몇 배로 늘어날 수 있습니다. 따라서 고성능 GPU 간 통신 라이브러리가 분산 학습의 성패를 결정하는 핵심 요소가 됩니다.

GPU 간 통신과 CPU 간 통신의 차이: NCCL과 MPI

분산 컴퓨팅에서 노드 간 데이터를 교환하는 통신 라이브러리는 워크로드 특성에 따라 다릅니다.

MPI(Message Passing Interface)는 여러 노드에 분산된 CPU 프로세스들이 서로 데이터를 주고받기 위한 표준 통신 인터페이스입니다. 엔지니어링 시뮬레이션에서 하나의 문제를 여러 노드에 나누어 계산할 때, 각 노드가 인접 노드의 계산 결과를 참조해야 하는데, 이때 MPI를 통해 노드 간 데이터를 교환합니다. MPI 사양을 구현한 라이브러리는 여러 가지가 있으며, 대표적으로 Open MPI와 Intel MPI가 있습니다. Open MPI는 오픈소스 커뮤니티에서 개발한 범용 구현으로 다양한 네트워크와 플랫폼을 지원하며, AWS EFA Installer에 기본으로 포함되어 있어 가장 쉽게 사용할 수 있습니다. Intel MPI는 Intel이 개발한 상용 구현으로, Intel 프로세서에 최적화된 성능 튜닝이 적용되어 있어 Intel CPU 기반 HPC 워크로드에서 강점을 가집니다.

NCCL은 GPU 간 통신을 위해 NVIDIA가 개발한 라이브러리입니다. 단일 노드 내 GPU 간 통신부터 다중 노드 간 GPU 통신까지 모두 지원하며, NVLink, NVSwitch, GPUDirect RDMA 등 GPU 전용 하드웨어 기능을 최대한 활용하도록 설계되어 있습니다. CPU가 네트워크 통신을 중재하는 전통적인 방식과 달리, NCCL은 GPU가 직접 다른 GPU와 통신할 수 있도록 하여 대규모 분산 학습에서 핵심적인 역할을 합니다.

AWS EFA 환경에서 이 두 라이브러리가 네트워크에 접근하는 경로도 다릅니다. Open MPI와 Intel MPI 같은 MPI 구현체는 libfabric API를 직접 사용하여 EFA 하드웨어에 접근합니다. NCCL은 libfabric을 직접 사용하지 않고, aws-ofi-nccl 플러그인을 통해 간접적으로 접근합니다. 이 플러그인이 NCCL의 통신 요청을 libfabric API 호출로 변환하는 브리지 역할을 합니다. 두 경우 모두 최종적으로는 EFA 하드웨어의 SRD(Scalable Reliable Datagram) 프로토콜을 통해 안정적인 데이터 전송이 이루어집니다.

설치 방식도 이 역할 차이에 따라 다릅니다. EFA Installer는 AWS에서 제공하는 통합 설치 스크립트로, EC2 인스턴스에서 EFA 네트워크를 사용하기 위해 필요한 소프트웨어를 한 번에 설치해줍니다. 여기에는 EFA 드라이버, libfabric, Open MPI, 그리고 aws-ofi-nccl 플러그인이 포함됩니다. NCCL 자체는 NVIDIA에서 별도로 제공하는 라이브러리이므로 따로 설치해야 합니다.

AWS에서 제공하는 전용 AMI를 사용하면 이러한 설치 과정을 생략할 수 있습니다. AWS Deep Learning AMI(DLAMI, GPU 버전) 에는 EFA Installer 구성 요소와 NCCL이 모두 사전 설치되어 있어 별도 설치가 필요 없습니다. AWS ParallelCluster AMI에도 EFA Installer 구성 요소가 사전 설치되어 있으며, 추가로 EFA Installer에는 포함되지 않는 Intel MPI도 환경 모듈 형태로 제공됩니다. ParallelCluster에서는 기본적으로 Open MPI가 경로에 설정되어 있으며, Intel MPI를 사용하려면 “module load intelmpi” 명령으로 전환해야 합니다. ParallelCluster를 사용하지 않는 환경에서는 Intel oneAPI HPC Toolkit으로 별도 설치해야 합니다.

EFA 소프트웨어 스택 구조

<그림1. EFA 소프트웨어 스택 구조>

위 그림처럼 MPI 라이브러리(Open MPI, Intel MPI)는 libfabric을 통해 네트워크 하드웨어와 통신합니다. libfabric은 일종의 네트워크 어댑터 모음으로, TCP·인피니밴드·EFA 등 다양한 네트워크를 지원하는 드라이버(provider)를 플러그인 형태로 가지고 있습니다. EFA를 사용하려면 libfabric에 EFA provider(그림1의 파란색 상자)가 포함되어 있어야 합니다.

여기서 주의할 점이 있습니다. EFA Installer가 설치하는 libfabric에는 EFA provider가 포함되어 있어 EFA 네트워크를 정상적으로 사용할 수 있습니다. 반면, Intel MPI가 자체 내장하고 있는 libfabric에는 EFA provider가 포함되어 있지 않아 EFA를 인식하지 못합니다. Intel MPI는 기본적으로 자체 내장 libfabric을 먼저 사용하기 때문에, 별도 설정 없이 실행하면 EFA 대신 TCP로 통신하게 되어 성능이 크게 저하됩니다. 이를 방지하려면 다음 환경변수를 설정하여 EFA Installer의 libfabric을 사용하도록 지정해야 합니다:

export I_MPI_OFI_LIBRARY_INTERNAL=0   # Intel MPI 내장 libfabric 사용 비활성화
export FI_PROVIDER=efa                # EFA provider 명시 지정

NCCL의 집합 통신이란?

분산 학습에서 여러 GPU가 협력하여 하나의 모델을 학습할 때, GPU들은 서로 데이터를 주고받아야 합니다. 이때 두 GPU 간의 일대일 통신을 반복하는 것은 수십~수백 개의 GPU가 참여하는 환경에서 매우 비효율적입니다. 집합 통신(Collective Communication)은 이 문제를 해결하기 위해 그룹 전체가 조율된 방식으로 동시에 데이터를 교환하는 통신 패턴이며, 앞서 소개한 NCCL은 GPU 간 집합 통신을 구현한 라이브러리입니다.

그림2의 6개 GPU 예시를 보면, 일대일 통신은 모든 쌍이 직접 연결되어야 해서 연결 수가 N × (N-1) / 2 = 15개입니다. 반면 링 구조에서는 각 GPU가 양옆 GPU와만 연결되므로 총 N = 6개의 링크면 충분합니다. GPU 개수 N이 커질수록 이 차이는 극적으로 벌어집니다. N=64면 일대일은 2,016개 연결이 필요하지만, 링은 64개 링크로 끝납니다. 약 31배 차이입니다. 연결 수가 적다는 것은 단순히 링크 개수의 문제가 아닙니다. 일대일 방식에서는 각 GPU가 동시에 여러 GPU로 데이터를 보내야 하므로 네트워크 대역폭이 여러 연결로 분할되어 경합이 발생합니다. 반면 링에서는 각 GPU가 이웃 1개와만 통신하므로 네트워크 대역폭을 최대로 활용할 수 있어, GPU 수가 늘어도 확장성이 유지됩니다.

일대일 통신 vs 집합 통신: GPU 간 데이터 교환 방식 비교

<그림2. 일대일 통신 vs 집합 통신: GPU 간 데이터 교환 방식 비교>

집합 통신을 자세히 들여다보려면 두 가지 층위를 구분하는 것이 중요합니다. 하나는 “무엇을 주고받느냐”, 즉 어떤 연산을 수행하는가(AllReduce, AllGather, ReduceScatter 등)입니다. 다른 하나는 “어떻게 주고받느냐”, 즉 GPU들이 어떤 경로로 데이터를 전달하는가(Ring, Tree 같은 알고리즘)입니다. 먼저 연산 종류부터 살펴보겠습니다.

NCCL의 집합 통신 연산

집합 통신 연산은 이름만 보면 낯설지만, Reduce / Gather / Scatter / All- 네 가지 개념의 조합으로 이해하면 직관적입니다.

Reduce

  • 여러 GPU의 값을 하나로 합치는 연산(합계, 평균, 최댓값 등)으로 개별 GPU가 계산한 그래디언트를 합산하는 것이 대표 예시입니다.

Gather

  • 여러 GPU에 흩어진 데이터를 합치지 않고 그대로 모아 하나의 큰 데이터로 이어붙이는 연산입니다.

Scatter

  • 하나의 데이터를 여러 GPU로 나누어 분배하는 연산입니다.

All- 접두사

  • 연산 결과를 특정 GPU 하나가 아니라, 참여하는 모든 GPU가 함께 공유한다는 의미입니다.

이 조합으로 연산 이름이 만들어집니다. AllReduce는 “Reduce(합산)한 결과를 All(모두)이 공유”, AllGather는 “Gather(모으기)한 결과를 All(모두)이 공유”, ReduceScatter는 “Reduce(합산)한 뒤 Scatter(나누어 분배)”한다는 뜻입니다. 키워드만 분해하면 어떤 연산인지 바로 그려집니다.

가장 빈번하게 사용되는 집합 통신 연산을 하나의 예시로 비교해보겠습니다. 3개의 GPU가 각각 3개 파라미터에 대한 그래디언트를 계산했다고 가정합니다. GPU 0은 [1, 2, 3], GPU 1은 [4, 5, 6], GPU 2는 [7, 8, 9]를 계산했습니다. 여기서 각 숫자는 해당 파라미터를 얼마나 조정해야 하는지를 나타내는 그래디언트 값입니다. 예를 들어, GPU 1은 파라미터 1을 4만큼, 파라미터 2를 5만큼, 파라미터 3을 6만큼 조정해야 한다고 계산한 것입니다.

AllReduce는 모든 GPU의 데이터를 파라미터별로 합산한 뒤, 그 결과를 모든 GPU에 동일하게 분배합니다. 위 예시에서 파라미터 1은 1+4+7=12, 파라미터 2는 2+5+8=15, 파라미터 3은 3+6+9=18로 합산됩니다. 연산이 끝나면 GPU 0, 1, 2 모두 [12, 15, 18]이라는 동일한 결과를 갖게 됩니다. 데이터 병렬 학습에서는 각 GPU가 서로 다른 데이터 배치로 계산한 그래디언트를 동기화할 때 사용됩니다. PyTorch DDP는 기본적으로 AllReduce로 합산한 뒤 각 GPU가 로컬에서 GPU 수로 나누어 평균을 구하며, 이 평균값이 파라미터 업데이트에 사용됩니다. AllReduce는 이렇게 매 학습 스텝마다 호출되므로, 분산 학습에서 가장 기본적이고 빈번하게 쓰이는 연산입니다. 참고로 NCCL AllReduce 자체는 SUM 외에도 AVG, MAX, MIN 등 여러 reduction 연산을 지원합니다.

AllGather는 합산 없이, 각 GPU의 값을 그대로 모아 모든 GPU에 제공하는 연산입니다. 앞서 본 예시에 AllGather를 적용하면 GPU 0, 1, 2 모두 [1, 2, 3, 4, 5, 6, 7, 8, 9]를 갖게 됩니다. 분산 학습에서는 주로 파라미터를 여러 GPU에 쪼개 저장하는 기법에서 사용됩니다. FSDP, ZeRO-3 같은 샤딩 데이터 병렬에서는 평상시 각 GPU가 전체 파라미터의 일부 조각만 보유합니다. 하지만 실제 forward/backward 연산이 실행되는 순간에는 그 레이어의 전체 파라미터가 필요하므로, AllGather로 흩어진 조각들을 모아 해당 레이어의 전체 파라미터를 각 GPU에 일시 복원합니다. 계산이 끝나면 즉시 해제하여 메모리 부담을 분산시킵니다.

ReduceScatter는 AllReduce처럼 파라미터별로 합산하지만, 결과를 모든 GPU에 복제하지 않고 쪼개서 나눠 가집니다. 합산 결과 [12, 15, 18] 중 GPU 0은 12만, GPU 1은 15만, GPU 2는 18만 보유합니다. 각 GPU가 전체 결과가 아닌 자기 몫만 들고 있으므로 메모리 사용량이 크게 줄어듭니다. 대표적으로 FSDP와 ZeRO-3의 역전파 단계에서 각 GPU가 계산한 그래디언트를 합산함과 동시에 자기 담당 파라미터 영역의 몫만 보관하는 데 사용됩니다.

정리하면, AllReduce는 “합산 후 전체를 모두에게”, AllGather는 “합산 없이 전체를 모두에게”, ReduceScatter는 “합산 후 나눠 가지기”로 이해하면 됩니다.

분산 학습에서는 매 학습 스텝마다 이러한 집합 통신이 호출됩니다. 기본 데이터 병렬(DDP)에서는 한 스텝당 AllReduce 한 번으로 끝나지만, 거대 모델 학습에서는 FSDP·ZeRO-3, Tensor Parallel, Pipeline Parallel 등 다양한 병렬화 기법이 조합되어 AllReduce, AllGather, ReduceScatter가 레이어마다 수차례 호출됩니다. 이 과정이 수만~수십만 번 반복되므로, 집합 통신의 효율이 전체 학습 시간을 좌우합니다.

NCCL의 자동 최적화

NCCL의 가장 강력한 기능 중 하나는 하드웨어 토폴로지를 자동으로 감지하고, 그에 맞춰 이러한 GPU 간 통신을 최적화하는 것입니다.

NCCL은 시스템 시작 시 GPU 간 연결 구조를 자동으로 파악합니다. NVSwitch의 존재 여부를 확인하여 모든 GPU가 연결된 구조인지 파악하고, 각 GPU가 NVLink를 통해 어떤 GPU와 직접 연결되어 있는지 매핑합니다. PCIe 계층 구조를 분석하여 CPU 소켓, PCIe 스위치, NUMA 노드 간의 관계도 이해합니다.

예를 들어, 듀얼 소켓 서버에서 GPU 0, 1, 2, 3이 CPU 0 소켓(NUMA 노드 0)에, GPU 4, 5, 6, 7이 CPU 1 소켓(NUMA 노드 1)에 연결되어 있다고 가정합니다. 같은 NUMA 노드 내 GPU끼리는 NVLink(NVLink가 없는 환경에서는 PCIe)를 통해 직접 통신하므로 빠르지만, 다른 NUMA 노드의 GPU와 통신하려면 CPU 간 연결을 거쳐야 하므로 상대적으로 느립니다. 이때 8개 GPU에서 AllReduce를 수행할 때, NUMA 구조를 무시하고 GPU 0과 GPU 5처럼 NUMA 노드를 넘나드는 통신을 빈번하게 배치하면 느린 경로를 반복적으로 타게 되어 전체 성능이 저하됩니다. NCCL은 이러한 토폴로지를 자동으로 감지하여 계층적으로 통신을 수행합니다. 먼저 NUMA 노드 0 안에서 GPU 0, 1, 2, 3끼리 합산하고, 동시에 NUMA 노드 1 안에서 GPU 4, 5, 6, 7끼리 합산한 뒤, 마지막에 두 그룹의 결과만 NUMA 노드 간에 교환합니다. 이렇게 하면 느린 경로를 타는 횟수를 최소화하여 집합 통신의 지연을 줄일 수 있습니다.

또한 EFA 카드와 같은 네트워크 인터페이스가 어떤 PCIe 스위치에 연결되어 있는지 파악하여, 각 GPU가 물리적으로 가장 가까운 NIC을 통해 데이터를 전송하도록 매핑합니다. 그림3과 같이 동일 PCIe 스위치에 있는 GPU와 NIC 조합은 PCIe 홉이 적어 레이턴시가 낮고, 다른 스위치를 거쳐야 하는 경우보다 빠릅니다. 이 정보를 바탕으로 GPU 간 최단 경로와 대역폭이 가장 높은 경로를 계산하여 통신 경로 맵을 생성합니다.

PCIe 토폴로지에 따른 GPU-NIC 매핑: 먼 경로 vs NCCL의 최단 경로 선택

<그림3. PCIe 토폴로지에 따른 GPU-NIC 매핑: 먼 경로 vs NCCL의 최단 경로 선택>

이렇게 감지된 토폴로지 정보를 기반으로 NCCL은 통신 파라미터를 자동으로 조정합니다. 링, 트리(Tree) 중에서 현재 하드웨어 구성에 가장 적합한 통신 알고리즘을 선택하고, GPU 개수와 NVLink 대역폭을 고려하여 병렬 통신 채널 수(동시에 동작하는 독립적인 통신 경로의 수)를 결정합니다. 예를 들어 H100처럼 NVLink 대역폭이 큰 GPU(GPU당 양방향 900 GB/s)에서는 단일 Ring만 쓰면 대역폭의 일부만 사용하게 되므로, 여러 개의 Ring을 동시에 구성하여 남는 대역폭까지 활용합니다.

또한 메시지 크기에 따라 전송 프로토콜을 자동으로 선택합니다. 여기서 메시지란 AllReduce 같은 통신 호출 한 번에 주고받는 전체 데이터를 뜻합니다. 작은 메시지에는 Low-Latency 프로토콜을 사용하여 데이터와 완료 신호를 함께 전송해 동기화 오버헤드를 줄이고, 큰 메시지에는 Simple 프로토콜을 사용하여 파이프라이닝(pipelining)으로 대역폭을 최대한 활용합니다. 알고리즘과 파이프라이닝의 상세 동작은 이어지는 섹션에서 다룹니다.

이러한 자동 튜닝 덕분에 개발자는 별도의 복잡한 설정 없이도 하드웨어의 최대 성능을 활용할 수 있습니다.

통신 알고리즘의 선택

NCCL은 데이터 크기와 GPU 개수에 따라 최적의 통신 알고리즘을 자동으로 선택합니다. 대표적인 알고리즘으로 앞서 언급한 링과 트리가 있습니다.

링 알고리즘은 GPU들이 원형으로 연결되어 데이터를 순차적으로 전달하는 방식입니다. 모든 GPU가 동시에 인접 GPU와 데이터를 주고받으므로 네트워크 대역폭을 최대한 활용할 수 있습니다. 다만 레이턴시 오버헤드가 GPU 수에 비례하여 커지므로, 수백 MB 이상의 대용량 데이터를 전송할 때 주로 사용됩니다.

트리 알고리즘은 계층적으로 데이터를 모으는 방식입니다. GPU 0과 GPU 1이 데이터를 합치고, GPU 2와 GPU 3이 합친 후, 이 두 결과를 다시 합치는 식으로 진행됩니다. 레이턴시가 O(log N)으로 링의 O(N)보다 낮지만, 각 단계에서 일부 GPU만 동시에 일하므로 대역폭 활용도는 링보다 낮습니다. 수 KB에서 수 MB 정도의 작은 데이터를 전송할 때 주로 사용됩니다.

예를 들어 AllReduce 연산에서 수 MB 이하의 작은 그래디언트는 트리 알고리즘이 유리합니다. 레이턴시가 O(log N)으로 링의 O(N)보다 낮아 레이턴시 측면에서 유리하기 때문입니다. 반면 수백 MB 이상의 대용량 파라미터는 링 알고리즘을 사용하여 대역폭을 최대한 활용합니다. 실제 소요 시간은 GPU 수와 네트워크 구성에 따라 크게 달라지므로, 정확한 수치는 nccl-tests로 환경에서 직접 측정하는 것이 정확합니다. NCCL은 이러한 알고리즘 선택을 데이터 크기와 GPU 개수를 기반으로 자동으로 수행합니다.

파이프라이닝과 청킹: 대용량 데이터를 잘게 나누어 동시에 전송하고 처리하기

링이든 트리이든 간에, NCCL은 대용량 데이터를 한 번에 보내지 않고 청크(chunk)라는 작은 조각으로 나누어 파이프라인 방식으로 처리합니다. 여기서 파이프라이닝은 하나의 통신 채널(예: 하나의 Ring) 내에서 일어나는 최적화이며, 앞서 설명한 여러 개의 병렬 채널과 독립적으로 함께 적용됩니다. 이해를 돕기 위해 링에서 인접한 두 GPU(송신자 GPU 0 → 수신자 GPU 1) 사이의 전송을 예로 들어보겠습니다. 실제 링에서는 이런 송수신 쌍이 원형으로 순차 연결되어 있고, 모든 쌍이 동시에 동작합니다.

청킹 없이 1GB 데이터를 전송한다면, GPU 0이 1GB 전체를 전송하는 동안 GPU 1은 아무것도 하지 못한 채 대기해야 합니다. 전송이 완료되어야 GPU 1이 수신한 데이터로 연산을 시작할 수 있습니다. 파일 하나를 통째로 다운로드한 후에야 열어볼 수 있는 것과 같습니다.

청킹과 파이프라이닝을 사용하면 상황이 달라집니다. 1GB 데이터는 수천 개의 청크로 나뉘고, GPU 0이 첫 번째 청크를 전송하자마자 GPU 1은 그것을 받아 자기 값과 합산하고 바로 다음 GPU로 전달합니다. 동시에 GPU 0은 이미 두 번째 청크를 전송하기 시작합니다. 이렇게 링의 여러 구간이 동시에 데이터를 나르는 상태가 되어 전송과 연산이 병렬로 이뤄지고, 전체 시간이 크게 단축됩니다. 동영상 스트리밍에서 전체 파일을 다운로드하지 않고도 재생이 시작되는 것과 같은 원리입니다.

청크 크기는 너무 작으면 청크 개수가 많아지면서 각 청크에 따라붙는 전송 헤더(송장 같은 관리 정보) 비중이 커져 실제 데이터 전송에 쓰이는 시간이 줄어듭니다. 반대로 청크가 너무 크면 첫 청크가 링을 한 바퀴 돌 때까지 오래 걸려서 뒷 GPU들이 그동안 놀게 되고, 파이프라인이 충분히 채워지지 못합니다.

여러 사람이 줄지어 서서 양동이로 물을 나르는 상황을 생각해봅시다. 아주 작은 컵으로 물을 나르면 사람 수는 많은데 매번 물 담고 넘기는 부수 동작이 많아 효율이 떨어집니다. 반대로 거대한 물통 하나로 한 번에 나르면, 첫 사람이 그 물통을 채워 다음 사람에게 건네는 동안 줄 뒤쪽 사람들은 계속 빈손으로 기다려야 합니다. 적당한 크기의 양동이로 빠르게 여러 번 넘기는 것이 전체 흐름을 끊지 않고 가장 빠릅니다.

NCCL의 청크 크기 결정도 같은 원리입니다. NCCL은 데이터 크기, 채널 수, GPU 개수, 네트워크 지연시간을 고려하여 최적의 청크 크기를 자동으로 결정하며, 노드 간 통신에서는 수백 KiB 수준의 청크를 파이프라인 단위로 사용합니다.

대규모 클러스터에서의 환경변수 설정은 어떻게 하면 될까?

NCCL은 기본 설정만으로도 하드웨어를 효율적으로 활용할 수 있도록 자동 최적화됩니다. 최신 소프트웨어 스택(libfabric 1 v2.5.1이상, aws-ofi-nccl v1.19.1 이상)을 사용하는 경우, 대부분의 환경변수는 기본값으로 충분하며 필수적으로 설정해야 하는 항목은 많지 않습니다.

현재 권장되는 최소 설정은 다음과 같습니다. FI_EFA_USE_HUGE_PAGE=0: PyTorch DataLoader에서 num_workers > 0로 멀티프로세스를 사용할 때 fork() 경로에서 huge page 할당 오류가 발생할 수 있습니다. 이를 회피하려면 명시적으로 0으로 설정합니다. 이전에는 FI_EFA_USE_DEVICE_RDMA=1(GPUDirect RDMA 활성화)이나 NCCL_PROTO=Simple 등의 설정이 필요했지만, 최신 버전에서는 자동으로 처리되므로 별도 설정이 불필요합니다.

디버깅과 성능 분석이 필요한 경우에는 NCCL_DEBUG=INFO를 설정하여 NCCL이 EFA를 정상적으로 사용하고 있는지, 어떤 알고리즘과 프로토콜이 선택되었는지 확인할 수 있습니다. 다만 이 설정은 대량의 로그를 출력하므로, 문제 진단 시에만 활성화하고 프로덕션 환경에서는 비활성화하는 것이 좋습니다.

앞서 설명한 채널 수도 비슷한 이유로 기본값을 권장합니다. 예를 들어 1GB 데이터를 전송할 때 채널이 4개면 각 채널이 250MB씩 담당하지만, 채널을 128개로 늘리면 각 채널에 약 8MB만 할당됩니다. 네트워크 전송에는 데이터 크기와 무관하게 발생하는 고정 오버헤드(패킷 헤더 처리, 전송 요청 설정 등)가 있기 때문에, 채널 수를 필요 이상으로 늘리면 각 채널에 할당되는 데이터가 너무 작아져 네트워크 대역폭을 충분히 활용하지 못하게 됩니다. 이러한 이유로 NCCL_MIN_NCHANNELS, NCCL_BUFFSIZE 등의 파라미터는 AWS의 awsome-distributed-training 가이드에서도 기본값 사용을 권장하고 있습니다. 앞서 설명한 파이프라이닝의 청크 크기도 NCCL 내부에서 데이터 크기, 채널 수, 네트워크 특성에 따라 자동으로 결정되며, 기본 자동 결정값이 대부분의 환경에서 잘 동작합니다.

실전에서 가장 중요한 원칙은 불필요한 튜닝에 시간을 쓰지 않는 것입니다. 값비싼 GPU 클러스터를 구축하고도 몇 주씩 nccl-tests만 반복하다가 실제 학습을 늦추는 경우가 자주 발생합니다. NCCL의 기본값은 AWS EFA 환경에서 대부분 최적에 가깝게 동작하도록 설계되어 있으므로, 대부분의 워크로드에서는 최소 설정(FI_EFA_USE_HUGE_PAGE=0)을 적용하고 NCCL_DEBUG=INFO로 EFA provider가 정상 인식되고 GPUDirect RDMA가 활성화되는지 로그로 확인한 뒤 바로 학습을 시작하면 됩니다. 기본값만으로도 EFA 성능의 대부분을 활용할 수 있습니다.

벤치마크와 환경변수 튜닝은 실제 학습에서 성능 병목이 드러난 이후에 진행하는 것이 효율적입니다. 이 경우에도 nccl-tests로 기준 성능을 먼저 측정한 뒤, 환경변수를 한 번에 하나씩만 변경하며 비교하는 것이 원칙입니다. 여러 변수를 동시에 바꾸면 어느 것이 효과를 냈는지 파악할 수 없습니다. 소프트웨어 버전별 상세 설정 프리셋과 인스턴스별 권장 구성은 awsome-distributed-training EFA Cheatsheet에서 확인할 수 있습니다.

다만 수천 GPU 이상의 대규모 클러스터에서는 기본값만으로 최적 성능을 달성하기 어려운 경우가 있습니다. AWS는 UltraCluster 아키텍처와 Capacity Block 또는 ODCR(On-Demand Capacity Reservation)을 통해 고성능 GPU 인스턴스를 동일한 네트워크 스파인 내에 배치하여 네트워크 경로를 최적화하지만, 워크로드의 메시지 크기 분포나 collective 패턴에 따라 NCCL 내부 파라미터를 조정하면 추가적인 성능 이점을 얻을 수 있습니다.

단일 노드 내 GPU 간 통신: NVSwitch 기반 아키텍처

GPU 간 통신은 크게 두 가지 범주로 나뉩니다. 단일 노드 내에서 GPU들이 직접 연결되어 통신하는 경우와, 여러 노드에 분산된 GPU들이 네트워크를 통해 통신하는 경우입니다. 각각의 통신 방식은 서로 다른 하드웨어 인프라와 소프트웨어 스택을 활용하며, 성능 특성도 크게 다릅니다.

단일 서버 내에서 GPU 간 통신은 NVLink와 NVSwitch를 통해 이루어집니다. 고성능 GPU 인스턴스인 P5와 P6는 8개의 H100 또는 B200/300 GPU를 탑재하고 있으며, 이들은 NVSwitch를 중심으로 완전한 메시 네트워크를 형성합니다. NVSwitch는 크로스바 스위치 역할을 수행하여 모든 GPU 간 직접 통신을 가능하게 합니다. 각 GPU는 NVLink를 통해 NVSwitch에 연결되며, NVSwitch는 모든 GPU 쌍 간에 full-bisection 대역폭을 제공합니다. 이는 어떤 GPU 쌍이 통신하더라도 동일한 최대 대역폭을 보장받는다는 의미입니다. 예를 들어, GPU 0과 GPU 7이 통신하는 동시에 GPU 2와 GPU 5가 통신하더라도, 두 통신 모두 최대 대역폭을 유지할 수 있습니다. 통신 경로는 “GPU → NVLink → NVSwitch → NVLink → GPU”의 형태를 띠며, 이 과정에서 CPU나 시스템 메모리를 거치지 않습니다. P5 인스턴스는 NVLink 4세대와 NVSwitch를 조합하여 각 GPU당 NVLink 양방향 900 GB/s의 대역폭을 제공합니다. P6 인스턴스는 2025년에 출시되었으며, NVIDIA B200/300 GPU와 NVLink 5세대를 탑재하여 P5 대비 2배 향상된 1,800 GB/s의 노드 내 GPU 통신 성능을 제공합니다. 이처럼 NVLink와 NVSwitch가 높은 대역폭과 낮은 지연시간을 제공하지만, 이러한 하드웨어 수준의 연결만으로는 충분하지 않으며, 효율적인 집합 통신 연산을 위해서는 NCCL이 필수적입니다.

멀티 노드 간 GPU 통신: EFA 기반 네트워크 스택

여러 노드에 분산된 GPU들이 통신할 때는 네트워크 인프라를 활용해야 합니다. AWS 환경에서는 EFA를 통해 고성능 노드 간 통신을 구현합니다. 최신 버전의 NCCL 사용을 권장하며, AWS Deep Learning AMI에는 검증된 최신 버전이 사전 설치되어 있습니다.

노드 간 GPU 통신 아키텍처

<그림4. 노드 간 GPU 통신 아키텍처>

노드 간 GPU 통신 아키텍처는 여러 계층의 소프트웨어와 하드웨어 구성 요소로 이루어져 있습니다. 애플리케이션 계층은 실제 딥러닝 프레임워크(PyTorch, TensorFlow 등)가 실행되는 영역입니다. 애플리케이션은 GPU 간 통신이 필요할 때 직접 네트워크 프로그래밍을 하지 않고, NCCL 라이브러리의 고수준 API를 호출하여 통신을 수행합니다.

NCCL 계층은 GPU 간 통신의 핵심 오케스트레이터 역할을 합니다. 애플리케이션으로부터 AllReduce, AllGather 같은 집합 통신 요청을 받으면, 이를 최적화된 통신 패턴으로 변환합니다. NCCL은 노드 내 GPU 간 통신(NVLink 활용)과 노드 간 GPU 통신(네트워크 활용)을 모두 관리하며, 하드웨어 토폴로지를 자동으로 감지하여 최적의 통신 경로를 선택합니다.

aws-ofi-nccl 플러그인 계층은 NCCL과 libfabric 사이의 브리지 역할을 합니다. NCCL은 원래 NVIDIA의 독자적인 통신 인터페이스를 사용하지만, AWS EFA 환경에서는 표준 libfabric API를 사용해야 합니다. aws-ofi-nccl 플러그인은 NCCL의 통신 요청을 libfabric API 호출로 변환하여, NCCL이 EFA를 투명하게 활용할 수 있도록 합니다. 이 플러그인이 없다면 NCCL은 일반 TCP/IP로 폴백(Fallback)하게 되어 성능이 크게 저하됩니다.

libfabric 계층은 앞서 언급한 것처럼, 고성능 네트워크 통신을 위한 표준 인터페이스를 제공하는 라이브러리입니다. 다양한 네트워크 하드웨어(인피니밴드, Ethernet, EFA 등)에 대한 추상화 계층을 제공하여, 상위 애플리케이션이 하드웨어의 세부 사항을 알 필요 없이 고성능 통신을 수행할 수 있게 합니다. libfabric은 EFA provider를 통해 EFA 하드웨어의 기능에 직접 접근합니다. 참고로 그림4에서 EFA provider는 libfabric 레이어에 포함되어 있습니다.

EC2 인스턴스가 부팅되면, Linux 커널이 EFA 드라이버(efa.ko)를 로드하여 EFA 하드웨어를 초기화하고, User space(커널 외부에서 실행되는 애플리케이션 영역)에서 하드웨어에 직접 접근할 수 있는 경로를 제공합니다. 이 경로를 통해 상위 계층인 libfabric이 커널을 거치지 않고 EFA 하드웨어와 직접 통신(OS-bypass)할 수 있습니다.

EFA 하드웨어는 AWS Nitro 시스템의 일부로, 각 EC2 인스턴스에 장착된 네트워크 카드입니다. 예를 들어 p5.48xlarge의 경우 32개의 독립적인 네트워크 카드가 존재하며, 각 카드는 100 Gbps의 대역폭을 제공합니다(총 3,200 Gbps). EFA 하드웨어는 AWS 자체 개발한 SRD(Scalable Reliable Datagram) 프로토콜을 하드웨어 수준에서 처리합니다.

p4d 이상 인스턴스에서는 EFA가 RDMA 연산을 지원하므로(p4d는 Read, p5 이상은 Read/Write), GPU와 네트워크 카드가 동일한 PCIe 스위치에 연결된 경우 GPUDirect RDMA를 활용할 수 있습니다. libfabric은 EFA 디바이스가 RDMA를 지원하는지 기능 속성(capability)을 확인하여 자동으로 판단하며, 최신 소프트웨어 스택에서는 별도 환경변수 설정 없이 GPUDirect RDMA가 자동으로 활성화됩니다. 이를 통해 데이터를 GPU 메모리에서 네트워크 카드로 직접 전송할 수 있어, CPU와 시스템 메모리를 거치는 오버헤드를 제거합니다.

두 GPU 인스턴스는 중앙의 네트워크 스위치 패브릭(Fabric)을 통해 연결됩니다. 이 패브릭은 AWS의 데이터센터 네트워크 인프라로, 수천 개의 인스턴스를 연결하는 대규모 스위칭 아키텍처입니다. AWS는 어떤 두 인스턴스 간에도 일관된 대역폭과 낮은 레이턴시를 보장합니다.

데이터 흐름 과정

실제 노드 간 GPU 통신이 발생하면 다음과 같은 흐름으로 데이터가 전달됩니다.

애플리케이션 → NCCL API 호출: 애플리케이션이 NCCL API를 호출하는 것으로 시작됩니다. 예를 들어, 분산 학습에서 PyTorch DDP는 gradient 동기화를 위해 AllReduce를 사용하도록 설계되어 있어, NCCL의 ncclAllReduce() 함수를 호출합니다. 즉, 어떤 집합 통신을 사용할지(AllReduce, AllGather 등)는 애플리케이션 계층이 결정합니다.
NCCL → 알고리즘 선택 및 청크 분할: NCCL은 요청받은 통신을 어떻게 실행할지를 최적화합니다. 초기화 시 하드웨어 토폴로지를 자동으로 탐지하여, 같은 노드 안의 GPU끼리는 NVLink로, 다른 노드의 GPU끼리는 네트워크(EFA)로 통신합니다. AllReduce 연산의 경우, NCCL은 데이터를 청크로 나누고 이 토폴로지 정보와 메시지 크기, GPU 규모를 바탕으로 링 알고리즘이나 트리 알고리즘 중 최적의 방식을 런타임에 선택합니다.
aws-ofi-nccl → libfabric API 변환: aws-ofi-nccl 플러그인이 NCCL의 통신 요청을 libfabric API 호출로 변환하여 다음 계층인 libfabric에 전달합니다.
libfabric EFA provider → OS-bypass 전송: libfabric의 EFA provider가 OS-bypass 경로로 EFA 하드웨어에 직접 접근하여 데이터를 전송합니다. GPUDirect RDMA가 활성화된 환경에서는 GPU 메모리에서 직접 데이터를 읽어 EFA 하드웨어로 전달하며, 이 과정에서 CPU와 시스템 메모리를 거치지 않습니다.
EFA 하드웨어 → SRD 패킷화 및 네트워크 전송: EFA 하드웨어가 데이터를 네트워크 패브릭으로 전송합니다. 데이터는 SRD(Scalable Reliable Datagram) 프로토콜로 패킷화되어 네트워크를 통해 원격 노드로 전달됩니다.
수신측 → GPUDirect RDMA로 GPU 메모리 직접 기록: 원격 노드에서는 EFA NIC가 수신한 패킷을 GPU 메모리에 직접 써넣습니다(GPUDirect RDMA). 수신 측 CPU 개입 없이(=원격 CPU 바이패스) 통신이 완료되므로 오버헤드가 최소화됩니다. 모든 데이터 전송이 완료되면 NCCL은 집합 통신 연산을 완료하고 애플리케이션에 제어를 반환합니다.

노드 내 vs 노드 간 통신 비교

노드 내 GPU 통신은 NVLink/NVSwitch를 통해 GPU당 양방향 900 GB/s(P5, NVLink 4세대) 또는 1,800 GB/s(P6, NVLink 5세대)의 대역폭을 제공합니다. 반면 노드 간 GPU 통신은 EFA를 통해 카드당 100 Gbps(p5.48xlarge 기준 총 3,200 Gbps)의 대역폭을 제공합니다. 노드 내 통신은 같은 서버 보드 위의 직접 연결인 반면, 노드 간 통신은 네트워크 스위치를 경유해야 하므로 대역폭과 레이턴시 모두에서 노드 내 통신에 비해 열위입니다. 그러나 대규모 분산 학습에서는 수백 개의 GPU를 활용해야 하므로, 노드 간 통신이 불가피합니다. NCCL의 소프트웨어 최적화와 EFA의 하드웨어 및 네트워크 최적화가 결합되어, AWS 환경에서도 온프레미스 인피니밴드 클러스터와 대등한 성능을 달성할 수 있으며, 대규모 분산 학습에서 수백 개의 GPU를 효율적으로 활용할 수 있습니다.

맺음말

지금까지 NCCL의 기본 개념부터 집합 통신 연산, 자동 최적화 메커니즘, 그리고 AWS EFA 환경에서의 노드 간 통신 아키텍처까지 살펴보았습니다.

이 블로그에서 다룬 내용을 정리하면 다음과 같습니다. NCCL은 단순한 통신 라이브러리가 아닙니다. 하드웨어 토폴로지를 자동으로 감지하고, 데이터 크기와 GPU 개수에 따라 최적의 알고리즘을 선택하며, 파이프라이닝과 청킹을 통해 계산과 통신을 오버랩시키는 지능적인 오케스트레이터입니다. 노드 내에서는 NVLink와 NVSwitch를 통해 초고속 저지연 통신을 처리하고, 노드 간에서는 EFA와 GPUDirect RDMA를 통해 CPU와 시스템 메모리를 거치지 않고 GPU 메모리에서 네트워크 카드로 직접 데이터를 전송합니다.

특히 AWS 환경에서 NCCL이 EFA를 제대로 활용하려면 aws-ofi-nccl 플러그인이 반드시 필요하다는 점을 기억해야 합니다. 이 플러그인이 NCCL의 통신 요청을 libfabric을 통해 EFA로 전달하는 연결 고리이기 때문입니다. 플러그인이 없으면 NCCL은 TCP/IP로 폴백하여 성능이 급격히 저하됩니다. 최신 소프트웨어 스택(libfabric 1 v2.5.1이상, aws-ofi-nccl v1.19.1 이상)에서는 GPUDirect RDMA가 자동으로 활성화되며, 대부분의 환경변수도 기본값으로 충분합니다. 대부분의 경우 기본값으로 바로 학습을 시작할 수 있으며, 튜닝은 실제 학습에서 성능 병목이 드러난 이후에 제한적으로 진행하는 것이 효율적입니다.

그러나 NCCL도 한계가 있습니다. AllReduce, AllGather와 같은 집합 통신은 모든 GPU가 동일한 크기의 데이터를 주고받는 대칭적 패턴에 최적화되어 있습니다. 그런데 최근 대규모 모델 학습에서 부상하는 MoE(Mixture-of-Experts) 모델은 입력 토큰에 따라 일부 전문가(expert) 네트워크만 선택적으로 활성화하는 구조입니다. 각 토큰이 어떤 전문가로 라우팅될지 매 스텝마다 달라지므로, GPU 간 전송량이 실행 시점에 결정되는 동적·비대칭 All-to-All 통신이 빈번하게 발생합니다. NCCL의 대칭적 집합 통신 모델로는 이런 불규칙 패턴을 효율적으로 처리하기 어렵습니다.

다음 블로그에서는 이 한계를 극복하는 GPU 중심 통신 패러다임을 다룹니다.

  • GPUDirect Async (IBGDA) — GPU가 CPU 개입 없이 직접 통신 명령을 발행
  • NVSHMEM — 여러 GPU의 메모리를 하나의 공유 주소 공간처럼 사용하는 SHMEM 프로그래밍 모델
  • PPLX-kernels — NVSHMEM 기반으로 MoE 워크로드를 AWS EFA 환경에서 최적화한 전용 커널

이 기술들이 어떻게 NCCL 중심 모델을 넘어서는지, 그리고 AWS EFA에서 어떻게 활용되는지 다음 글에서 자세히 살펴보겠습니다.

Sangman Cho

Sangman Cho

조상만 Solutions Architect는 AWS 입사 이후, Automotive 및 Manufacturing 고객의 클라우드 기반의 디지털 전환 업무를 지원하였으며, 현재는 AWS 코리아 전체의 고성능 컴퓨팅(HPC)과 양자 컴퓨팅 등 계산 과학 영역의 디지털 전환 업무를 지원하고 있습니다.