CUDA Multicast

👉 CUDA multicast는 하나의 데이터를 여러 GPU에 동시에 전달해서 GPU 간 통신을 효율화하는 기술입니다.

 

구분	Unicast	Broadcast	Multicast
통신 방식	1 → 1	1 → 전체	1 → 특정 그룹
대상	한 개	모든 노드	선택된 여러 노드
효율성	낮음 (여러 번 전송 필요)	비효율적 (불필요한 대상 포함)	가장 효율적
네트워크 부담	높음	매우 높음	낮음
example	한 사람에게 전화 * GPU0 → GPU1	마이크로 모든 사람에게 방송 * GPU0 → 모든 GPU	특정 그룹방에만 메시지 전송 * GPU0 → GPU1, GPU3, GPU5

 

구현 방식

기본 NCCL ↔ CUDA 관계 구조

[사용자 코드]
     ↓
[NCCL (고수준 통신)]                   

     * CUDA 기반으로, multi-GPU 통신 패턴 (broadcast, all-reduce, ...) 제공.

     * 각 패턴 내부에서, hardware topology (NVLink/NVSwitch 여부, PCIe tree 구조, NUMA 등)를 보고, 어떤 알고리즘이 적합할지 (ring? tree? NVLS? ...), 그래서 data를 어떤 GPU ↔ GPU 경로로 보낼지 (direct? relay?) 결정 최적화.

     ↓
[CUDA (저수준 실행 + 메모리)]   

     * low-level에서 GPU 메모리 관리 / 커널 실행 / GPU ↔ GPU 데이터 복사
     ↓
[GPU HW (NVLink, PCIe, NVSwitch)]

 

Unicast

NCCL P2P Send/Recv:

여기서는 P2P 통신이기 때문에 ring/tree 같은 collective 알고리즘까지 고려하지는 않는다.

다만 hardware topology 고려해서 NVLink (또는 NVLink 없으면 PCIe P2P)로 direct GPU ↔ GPU transport할지 / 아니면 아예 P2P 불가능하다면 Host staging ( GPU ↔ CPU ↔ GPU) 할지 정도는 고려.

* `NCCL_P2P_DISABLE=1`: NVLink나 PCIe를 사용하는 direct GPU-to-GPU P2P transport를 끄는 옵션.

* `NCCL_P2P_LEVEL`: 

// rank 0 -> rank 1
if (rank == 0) {
  ncclSend(sendbuf, count, ncclFloat, 1, comm, stream);
} else if (rank == 1) {
  ncclRecv(recvbuf, count, ncclFloat, 0, comm, stream);
}

 

NCCL 없이 CUDA만으로 P2P Copy 구현:

// GPU0 -> GPU1 로 1회 전송
cudaSetDevice(0);
float* src;
size_t bytes = N * sizeof(float);
cudaMalloc(&src, bytes);
// src 채우기 ...

cudaSetDevice(1);
float* dst;
cudaMalloc(&dst, bytes);

// device 0의 메모리를 device 1로 복사
cudaMemcpyPeerAsync(dst, 1, src, 0, bytes, stream1);

Broadcast

NCCL Broadcast:

NCCL이 ring/tree/NVLS 등 사용 가능한 알고리즘 중에서 topology에 맞춰 선택하므로, root 단일 병목이 줄고 전체 fabric을 더 잘 쓴다. 공식 문서상 기본값은 topology와 architecture에 따라 자동 선택.

// 모든 rank가 같은 collective 호출
ncclBroadcast(
    sendbuff,     // root에서만 의미 있음
    recvbuff,
    count,
    ncclFloat,
    /*root=*/0,
    comm,
    stream);

Multicast

Multicast는 Broadcast의 좀더 general한 개념. broadcast는 "모든" GPU에게 방송하는 방식인 반면, multicast는 "여러 몇몇" GPU에게 방송하는 방식.

NCCL에는 “사용자가 직접 `ncclMulticast(...)`를 호출하는 API”는 없다.

앱 코드에서는 여전히 `ncclBroadcast(...);` 또는 `ncclAllReduce(...);` 호출. 

그런데 내부적으로 지원 환경이면 multicast/offload 성격의 hardware 경로를 사용할 수 있다. 

대신 NCCL 문서에 따르면 알고리즘 선택 변수 `NCCL_ALGO`에 `NVLS`, `NVLSTree`가 있으며, 이들은 NVLink SHARP offload를 사용.

• 최고 성능 가능성: multicast-capable fabric일 때
• 하지만 조건이 까다로움: NVSwitch/NVLink SHARP 지원, 라이브러리/알고리즘 지원 필요
• 일반 PCIe-only 환경에서는 이런 이점이 없습니다

Q. NVLink SHARP?
A. GPU 간 통신 + 연산을 “네트워크(NVSwitch)” 안에서 처리하는 기술.
SHARP = Scalable Hierarchical Aggregation and Reduction Protocol. 네트워크 switch 안에서 reduce (sum, max 등) 연산을 수행하게 하는 기술이다. 원래는 노드 간 통신 때 쓰이는 InfiniBand에서 나온 기술임.

이걸 NVIDIA intra-node 환경(NVLink + NVSwitch)으로 끌고 와서, NVLink SHARP란, 이 SHARP 기술을 같은 노드 내 GPU 간 네트워크 (NVLink + NVSwitch)에 적용한 것이다. 따라서 NVSwitch가 단순히 네트워크 switch가 아니라 data 복제 + reduction 연산까지 수행할 수 있게 한다. (GPU ↔ NVSwitch ↔ GPU)

Ex) 기존 software 기반 allreduce는 GPU가 직접 연산을 하고 통신량이 많다:
     GPU0 → GPU1 → GPU2 → GPU3 (데이터 이동 + 각 GPU에서 sum)

반면 NVLink SHARP는 각 GPU가 NVSwitch로 데이터를 보내면 NVSwitch가 연산하고 결과를 각 GPU에 분배한다:
              NVSwitch
             /       |       \
GPU0 GPU1 GPU2 GPU3

즉, NVLink SHARP offload란 말은, 원래 GPU가 하던 연산을 NVSwitch한테 offloading한다는 뜻!

 

DGX A100 Hardware Topology

참고 사항) 아래 그림은 NVIDIA DGX A100 노드의 hardware topology이다. 하단부를 보면 GPU 8대가 각각 모두 6대의 NVSwitch에 NVLink로 연결되어 있고, NVSwitch가 중간다리 역할을 해주고 있다. 

그리고 다른 DGX 노드에 data를 전송하기 위해서는:

GPU memory -> PCIe Bus -> PCIe Switch -> PCIe Bus -> NIC (InfiniBand HCA) 
    -> [InfiniBand Network: InfiniBand Cable -> InfiniBand Switch(es) -> InfiniBand Cable]
        -> Destination Node NIC -> ... -> Destination GPU memory

이렇게, NVSwitch 말고 PCIe Switch를 거쳐서 InfiniBand로 연결된다. (CPU도 안 거친다. a.k.a. GPU direct RDMA)

NVIDIA DGX A100 host (official 8-card machine) hardware topology 출처: https://www.naddod.com/blog/high-performance-gpu-server-hardware-topology-and-cluster-networking-2?srsltid=AfmBOooe4aSDdrvDSFdmrJDCWpPYn7Fr81pEGr6cTsLqH3s9gK1JiHkH