NCCL and NVSHMEM — two GPU communication models, side by side

PyTorch 에서 dist.all_reduce 한 줄을 부르면 그게 NCCL 로 들어간다. 학습 코드의 99% 가 그 한 줄로 끝난다. 그런데 frontier inference / training 의 최근 코드 (DeepSeek, FlashInfer, NVSHMEM-based MoE) 는 NVSHMEM 을 직접 쓴다. 같은 NVLink, 같은 RDMA 위에서 — 왜 다른 인터페이스인가?

강의의 출발 질문 셋.

1. NCCL 의 collective 추상이 왜 부족할 때가 있는가 — host-launched, batched, bulk 동기 모델 자체의 한계.
2. NVSHMEM 의 PGAS 모델은 어떤 패턴을 추가로 표현 가능하게 하는가 — fine-grained, GPU-initiated, asynchronous put/get.
3. 두 모델은 같이 쓰는가, 둘 중 하나만 쓰는가 — 답: 같이 쓴다. layer 가 다르다.

Jeff 의 입장은 “두 모델은 경쟁이 아니라 보완”이다. NCCL 은 큰 단위의 collective를 가장 빠르게, NVSHMEM 은 작은 단위의 비대칭 통신을 가장 자연스럽게. 이게 강의 전체를 끌고 가는 framing.

NCCL 의 핵심 단어는 collective. 모든 rank 가 같은 함수를 같은 인자로 부르고, 모두가 끝나야 다음 코드가 실행된다. ncclAllReduce, ncclBroadcast, ncclAllGather, ncclReduceScatter — 이 4-5개가 학습 코드의 99%.

// NCCL 의 표준 사용 패턴 — host code
ncclComm_t comm;
ncclCommInitRank(&comm, nranks, id, rank);

// 각 rank 가 자기 GPU 에 sendbuff/recvbuff 준비
cudaMalloc(&sendbuff, size);
cudaMalloc(&recvbuff, size);

// 모든 rank 가 같이 부른다
ncclAllReduce(sendbuff, recvbuff,
              count, ncclFloat,
              ncclSum, comm, stream);

// stream 위에서 비동기로 실행됨 — 이후 sync 필요
cudaStreamSynchronize(stream);

PGAS (Partitioned Global Address Space) 의 약속 — 각 GPU 의 메모리가 모두 같은 가상 주소 공간 안에 있다. 단, 물리적으로는 분산. 각 메모리 영역은 한 GPU 에 속한다 (partitioned). 다른 GPU 에 있는 변수에 접근하려면 shmem_get, 다른 GPU 에 쓰려면 shmem_put. 한쪽만 부르면 끝난다 — one-sided.

// NVSHMEM 의 표준 host code
nvshmem_init();
int me = nvshmem_my_pe();
int npes = nvshmem_n_pes();

// 모든 PE 가 같은 사이즈, 같은 offset
float* buf = (float*)
    nvshmem_malloc(N * sizeof(float));

// PE 0 이 PE 1 의 buf[0..N) 를 자기 src 로 덮어씀
if (me == 0) {
    nvshmem_float_put(buf, src, N, 1);
}
// PE 1 은 아무것도 안 부른다 — 자기 buf 가 바뀐다
nvshmem_barrier_all();

nvshmem_free(buf);
nvshmem_finalize();

// Device 에서 직접 — kernel 안에서 통신
__global__ void exchange(float* buf) {
    int tid = threadIdx.x;
    int me  = nvshmem_my_pe();
    int peer = (me + 1) % nvshmem_n_pes();
    // thread 단위 fine-grained put
    nvshmem_float_p(&buf[tid], src[tid], peer);
}

왜 GPU-initiated 가 결정적인가

NVSHMEM 의 진짜 강점은 device-side API. 위 두 번째 코드처럼 kernel 안에서 직접 nvshmem_float_p 를 부른다 — kernel 을 끝내고 host 로 나갔다가 NCCL collective 를 부르는 게 아니라, 같은 kernel 안에서 통신과 계산이 섞인다.

두 라이브러리를 한 표 위에 펼쳐놓는다. 같은 hardware 위에서 다른 약속을 한다는 점이 명확해야 한다.

NCCL 의 black box 안에 들어가면 같은 allreduce 가 ring 알고리즘과 tree 알고리즘 두 형태로 다르게 구현돼 있다. 작은 메시지에서는 tree (latency 우위), 큰 메시지에서는 ring (bandwidth 우위). NCCL 이 자동으로 고른다.

double binary tree

현대 NCCL 의 기본 — 두 개의 보완적 binary tree 를 동시에 돌려 tree 의 bandwidth 약점을 보강한다. 각 tree 가 서로 다른 link 만 사용하도록 짜여 있어서 NVLink 의 양방향 대역폭을 모두 활용. 이게 NCCL 2.x 가 “fastest collective library” 의 자리에 머물게 한 핵심 알고리즘.

NCCL 은 host 가 launch 하는 collective. 그 결과 — kernel 이 끝나야 통신이 시작된다 (또는 별도 stream 위에 띄워야 함). NVSHMEM 은 kernel 안에서 직접 put/get. 한 thread block 이 작업을 끝내자마자 그 결과를 다른 GPU 로 send 시작.

실전 예 — DeepEP / DeepSeek 의 expert parallelism dispatch. token 마다 어느 expert 로 갈지가 다르다. 이걸 NCCL 의 모든 token batched all-to-all 로 보내면 padding 과 sync 비용이 크다. NVSHMEM 의 per-token put 으로 보내면 정확히 필요한 만큼만, kernel 안에서 직접.

NCCL 과 NVSHMEM 모두 tightly coupled 모델. communicator 안의 한 GPU/노드가 죽으면 collective 전체가 멈추고, 보통 process group 을 다시 만들어야 한다. 큰 학습 클러스터 (10k GPU) 에서는 이게 사실상 항상 일어난다.

communication 의 cost 모델 — T(n) = α + β·n. α 는 latency floor, β 는 reciprocal bandwidth, n 은 메시지 사이즈. 이 단순한 모델 위에서 알고리즘 선택의 모든 기준이 나온다.

그리고 — “latency 와 bandwidth 는 같은 hardware 위에서도 다른 측정치”. NVLink 의 peer bandwidth 는 ~600 GB/s 인데 latency 는 ~1 µs 영역. 같은 NVLink 가 어떤 메시지에 대해서는 빠르고, 어떤 메시지에 대해서는 어쩔 수 없이 늦다 — 사이즈가 작으면 α 가 지배.

강의의 끝부분은 “NCCL 이 어디로 가는가”. 흥미롭게도 NCCL 의 진화 방향이 NVSHMEM 의 기능을 흡수하는 방향이다 — 같은 라이브러리 안에서 두 모델을 모두 제공.

• NCCL symmetric memory (NCCL 2.20+) — NVSHMEM 의 symmetric heap 을 NCCL 에 도입. PyTorch 의 SymmetricMemory API 가 이걸 노출.
• NCCL device API — kernel 안에서 직접 부를 수 있는 NCCL primitive 일부. 작은 메시지 case 를 위한 fast path.
• fault tolerance — recoverable communicator. 일부 rank 가 사라져도 남은 rank 끼리 collective 를 계속할 수 있도록 (확인 필요, NCCL 3.x 의 design 단계).
• topology-aware tree — multi-rail IB 와 NVLink 5 의 새 topology 에 맞춘 알고리즘.

손에 새기기 — 실습 시퀀스

1. nccl-tests bandwidth sweep — ./all_reduce_perf -b 8 -e 1G -f 2 -g 4. 8B 부터 1GB 까지 sweep. α 와 β 를 그래프로 분리.
2. NCCL_DEBUG=INFO — 학습 코드 한 번 돌려 NCCL 이 ring vs tree 중 무엇을 골랐는지 직접 확인. NCCL INFO Channel/Tree 라인.
3. NVSHMEM hello world — ParRes/Kernels 의 stencil 예제. 두 PE 사이의 halo exchange 를 NVSHMEM put 으로 직접 구현.
4. device-initiated put — kernel 안에서 nvshmem_float_p 를 부르는 커널 작성. host-launched NCCL allreduce 와 latency 비교.
5. PyTorch SymmetricMemory — torch 2.4+ 의 새 API. NCCL symmetric memory 위에서 작은 메시지 communication 패턴 직접 작성.
6. NCCL_ALGO 강제 — 같은 allreduce 를 NCCL_ALGO=Ring 과 NCCL_ALGO=Tree 로 두 번 돌려 메시지 사이즈별 차이 측정. crossover 직접 찾기.

이 강의의 두 도구가 다른 강의들에서 어떻게 다시 등장하는지 — communication 가족의 핵심 노드.

• NCCL 3.x 의 정확한 fault tolerance design — 본문에서 “검토중” 으로 적은 부분. 공식 NCCL 릴리스 노트와 NVIDIA GTC 자료 확인 필요.
• ring vs tree crossover 의 정확한 사이즈 — H100 NVLink, A100 NVLink, B200 NVLink 별로 다름. nccl-tests 자체 측정 필요.
• NVSHMEM 와 IBGDA 관계 — IB GPUDirect Async 의 상태. NVSHMEM 이 사용하는 transport. 강의에서 잠깐 언급됐는지 영상 재확인 필요.
• PyTorch SymmetricMemory 의 실 사용 사례 — 어떤 production 코드가 이미 채택했는지. 본 노트에서는 일반화된 표현으로만.
• α-β 모델의 한계 — 단순 모델은 contention, congestion, NUMA 효과를 못 잡음. 실제 측정에서는 변동 큼.