GPU 분산 훈련, EFA와 InfiniBand의 선택? DeepEP vs PPLX-kernels
분산 훈련(Distributed Training)에서 GPU 간 통신 성능이 전체 훈련 효율에 미치는 영향과 데이터 경로(Data Path) 및 제어 경로(Control Path) 관점에서 기술 발전 과정을 설명
GPUDirect RDMA를 통해 CPU 개입을 줄이고, EFA(Elastic Fabric Adapter)와 결합하여 OS 네트워크 스택까지 우회하는 직접 전송 경로를 구현
DeepEP는 인피니밴드(InfiniBand) 환경에서 IBGDA(InfiniBand GPUDirect Async)를 활용하여 낮은 레이턴시를 달성하지만, AWS EFA 환경에서는 성능 저하 발생
PPLX-kernels는 EFA 환경에 맞춰 GDRCopy를 활용, CPU 프록시(CPU Proxy)와 GPU 간 동기화 오버헤드를 줄여 DeepEP보다 우수한 성능을 보임
하드웨어 스펙보다 네트워크 인터커넥트(Network Interconnect)의 설계 사상을 이해하고, 환경에 맞는 최적화를 적용하는 것이 중요함을 강조
GPU 간 통신 기술 발전 과정: 데이터 경로 vs. 제어 경로
본문은 분산 훈련에서 GPU 간 통신 성능 향상을 위해 CPU 개입을 줄이는 방향으로 기술이 발전해 왔음을 설명한다.
GPUDirect RDMA(GDR): GPU 메모리에서 CPU 메모리를 거치지 않고, NIC(Network Interface Card)이 PCIe를 통해 GPU 메모리에 직접 접근하여 데이터 전송. CPU 메모리 복사 제거를 통해 레이턴시 감소
GPUDirect Async(IBGDA): 전송 명령(Work Request)을 NIC에 직접 전달하여 제어 경로(Control Path)의 CPU 개입 제거. NVIDIA 벤치마크에 따르면, 1KB 미만 소규모 메시지에서 기존 대비 최대 9.5배 처리량 향상
결과적으로, 데이터 경로와 제어 경로 모두에서 CPU 개입을 제거하여 GPU와 NIC 간 직접 통신을 실현. 이는 소규모 메시지 빈번한 통신(Small Message Frequent Communication) 환경에서 특히 중요하며, MoE 모델과 같은 워크로드에 효과적이다.
AWS EFA와 InfiniBand의 설계 철학 비교
AWS EFA와 InfiniBand는 각기 다른 설계 철학을 바탕으로 분산 훈련 환경을 구축한다. 인터커넥트(Interconnect) 기술 선택은 하드웨어 스펙보다 환경에 맞는 최적화가 중요함을 시사한다.
InfiniBand: NVIDIA가 하드웨어(ConnectX NIC)부터 소프트웨어(NVSHMEM, NCCL, IBGDA)까지 수직 통합(Vertical Integration)된 생태계. 자사 기술 스택에 최적화된 워크로드에서 뛰어난 성능 발휘
AWS EFA: SRD(Scalable Reliable Datagram) 프로토콜 기반의 멀티패스 전송, OS-바이패스, 수천 노드 규모의 확장성을 목표로 설계된 클라우드 네이티브(Cloud Native) 인터커넥트. 유연성과 확장성에 초점
DeepEP와 PPLX-kernels 사례: 인피니밴드 전용 기능(IBGDA)에 강하게 결합된 DeepEP는 EFA 환경에서 성능 저하. PPLX-kernels는 EFA 특성에 맞게 GDRCopy, Unified Memory, 벌크 RDMA write 등을 활용하여 EFA 환경에 최적화된 성능을 달성.
MoE 모델 최적화: DeepEP vs. PPLX-kernels
MoE(Mixture-of-Experts) 모델의 Expert Parallelism 통신 최적화는 하드웨어 환경에 따라 다른 접근 방식을 취한다. 각 환경에 맞는 최적화 전략이 필요하다.
DeepEP: 인피니밴드 + IBGDA에 최적화된 솔루션. NVSHMEM을 기반으로, GPU가 CPU 개입 없이 NIC에 직접 전송 명령을 내리는 방식으로 극도로 낮은 레이턴시(Low Latency) 달성
PPLX-kernels: AWS EFA 환경에 최적화된 솔루션. CPU 프록시와 GPU 간 동기화 병목을 줄이기 위해 GDRCopy 활용. GDRCopy를 통해 CPU-GPU 간 동기화 레이턴시를 최소화하여 IBGDA 부재의 성능 격차를 메움
Perplexity AI의 PPLX-kernels는 ConnectX-7 환경의 DeepEP를 포함한 모든 IBGDA 기반 구현을 상회하는 성능을 달성. 이는 하드웨어 스펙보다 인프라 특성을 이해하고 최적화하는 것이 중요함을 보여준다.
GDRCopy: CPU-GPU 동기화 병목 해결
GDRCopy는 CPU 프록시 기반 구조에서 발생하는 동기화 병목을 해결하기 위한 핵심 기술이다. CPU와 GPU 간 데이터 전송 효율성 향상에 기여한다.
문제점: CPU가 GPU 메모리를 읽기 위해 cudaMemcpy 함수 사용. 이는 수백 MB 대용량 데이터 전송에 최적화되어 있어, 수 바이트 신호 값 하나를 읽는 데도 약 7μs 대기 시간 발생
GDRCopy 해결책: CPU가 GPU 메모리를 자신의 메모리처럼 직접 읽을 수 있도록 미리 매핑. 복사 엔진을 거치지 않고 CPU가 GPU 메모리에 직접 접근 가능
결과: 신호 하나를 읽는 데 걸리는 시간 1μs 미만으로 감소. PPLX-kernels에서 CPU 프록시와 GPU 커널 사이의 동기화를 빠르게 처리하기 위해 활용. MoE 워크로드의 성능 향상에 기여
NVSHMEM: MoE 모델의 불균등 통신 최적화
NVSHMEM은 MoE 모델의 Expert Parallelism에서 발생하는 불균등한 AlltoAll 통신을 효율적으로 처리하기 위한 기술이다. PGAS(Partitioned Global Address Space) 모델을 활용하여 GPU 간 직접 통신을 가능하게 한다.
NCCL의 한계: AllReduce, AllGather와 같은 집합 통신에 최적화. MoE의 불균등한 라우팅 패턴에서는 패딩(Padding)으로 인한 대역폭 낭비 발생
NVSHMEM의 장점: PGAS 모델을 사용하여 원격 GPU 메모리를 로컬 메모리처럼 주소로 직접 참조. GPU가 연산 도중에 nvshmem_put()을 호출하여 CPU를 거치지 않고 즉시 원격 GPU 메모리에 데이터 쓰기 가능
결과: GPU가 연산 결과가 나오는 즉시 통신 시작. MoE의 불균등한 라우팅 패턴에서 효율적인 통신을 가능하게 하며, CPU 병목을 줄이고 패딩 없이 실제 데이터만큼 전송 가능
