이질 클러스터 — MPI + CUDA

20.1 배경 — 노드를 넘는 순간 게임이 바뀐다

여기까지 우리는 한 GPU 안의 세계를 다뤘다. 한 GPU가 좀 부족하면 한 노드에 GPU를 4~8장 꽂는다. 그래도 부족하면? 노드를 여러 대 묶어 클러스터로 만든다. HPC(High Performance Computing) 클러스터의 흔한 모습이다. 노드 하나에 CPU + GPU 4~8장 + InfiniBand NIC. 노드 사이는 InfiniBand 또는 RoCE 같은 고속 네트워크로 연결된다.

단일 GPU에서 무시해도 됐던 한 가지가 갑자기 dominant 비용이 된다 — 노드 간 통신이다. NVLink로 연결된 같은 노드 내 GPU 사이는 수백 GB/s, 수 마이크로초 latency. InfiniBand 노드 간 통신은 수십 GB/s, 수 마이크로초 latency지만 호스트 메모리 경유 시 수십 마이크로초까지 늘어난다. 계산은 GPU에서 빠른데, 통신이 느려 결국 통신에 지배된다는 풍경이 흔하다.

클러스터 환경에서 사실상 표준 통신 라이브러리가 MPI(Message Passing Interface)다. 1990년대부터 표준화되어 지금도 쓰이는 이유는 단순하다 — 명시적 메시지 전달 모델이 분산 메모리 환경에 가장 솔직하게 맞는다. 누가 어떤 데이터를 가지고 있고, 누가 누구에게 보내야 하는지를 코드로 직접 적는다. 시작은 까다롭지만, 일단 익히면 어떤 분산 알고리즘이든 표현할 수 있다.

20.2 러닝 예제 — 8장의 스텐실을 클러스터로

8장에서 다룬 3D 스텐실(stencil) 계산을 확장해 보자. 큰 격자 N_x×N_y×N_z에서 매 스텝 각 점이 자기와 이웃의 값을 합산해 갱신한다. 격자 전체가 한 GPU에 안 들어간다고 가정하자. 어떻게 자를 것인가?

가장 자연스러운 것이 도메인 분할(domain decomposition)이다. 격자를 z 방향으로 P개 슬랩(slab)으로 나누고, 슬랩 하나당 한 노드(또는 한 GPU)에 할당한다. 슬랩 i는 자기 슬랩 안의 점들을 갱신할 수 있다. 다만 슬랩의 위/아래 경계점은 인접 슬랩의 값이 필요하다 — 여기서 통신이 발생한다.

       Rank 0          Rank 1          Rank 2          Rank 3
   ┌───────────┐  ┌───────────┐  ┌───────────┐  ┌───────────┐
   │           │  │           │  │           │  │           │
   │  slab 0   │  │  slab 1   │  │  slab 2   │  │  slab 3   │
   │           │  │           │  │           │  │           │
   │ ── halo ──│←→│── halo ───│←→│── halo ───│←→│── halo ── │
   └───────────┘  └───────────┘  └───────────┘  └───────────┘
                  ↑           ↑
              boundary 면을 매 스텝 교환

그림 20.1 — 슬랩 분할과 halo. 각 rank는 자기 슬랩 + 위/아래 halo 한 층을 들고 있다.

각 rank는 자기 슬랩 + 그 위아래의 한 줄(halo)을 메모리에 가지고 있다. 매 스텝마다 (a) halo를 이웃과 교환, (b) 자기 슬랩 갱신. 이 단순한 패턴을 클러스터 위에서 맞물려 돌리는 것이 우리 목표다.

20.3 MPI 기초 — rank, communicator, 점대점 통신

MPI 프로그램은 동일한 바이너리를 P개 프로세스로 실행한다(보통 mpirun -np P ./app). 각 프로세스는 0..P-1의 rank를 받는다. rank는 자기 신원증이다. communicator는 프로세스 그룹의 핸들 — 가장 흔한 것이 모든 프로세스를 포함하는 MPI_COMM_WORLD.

#include <mpi.h>
#include <cuda_runtime.h>

int main(int argc, char** argv) {
    MPI_Init(&argc, &argv);

    int rank, size;
    MPI_Comm_rank(MPI_COMM_WORLD, &rank);
    MPI_Comm_size(MPI_COMM_WORLD, &size);

    // 한 노드에 GPU 여러 장이면 rank 기반으로 골라잡기
    int devs; cudaGetDeviceCount(&devs);
    cudaSetDevice(rank % devs);

    // ... 본 작업 ...

    MPI_Finalize();
    return 0;
}

가장 단순한 통신은 점대점(point-to-point)이다. MPI_Send로 보내고 MPI_Recv로 받는다.

// rank 0 -> rank 1 로 N개 float 보내기
if (rank == 0) {
    MPI_Send(buf, N, MPI_FLOAT, /*dest*/1, /*tag*/42, MPI_COMM_WORLD);
} else if (rank == 1) {
    MPI_Status st;
    MPI_Recv(buf, N, MPI_FLOAT, /*src*/0, /*tag*/42, MPI_COMM_WORLD, &st);
}

MPI_Send는 라이브러리/구현에 따라 작은 메시지는 즉시 반환하지만(eager), 큰 메시지는 받는 쪽이 받기 전까지 블록될 수 있다(rendezvous). 이 모호함 때문에 deadlock이 자주 난다 — A는 B에게 보내려 블록, B도 A에게 보내려 블록, 둘 다 받기 안 시작했으니 영원히 정지. 그래서 실전에서는 비블로킹 버전을 쓴다.

20.4 비블로킹 통신과 halo exchange 패턴

비블로킹(nonblocking) 점대점은 MPI_Isend, MPI_Irecv이다. 이름에 I(immediate)가 붙는다. 둘 다 MPI_Request를 돌려주고, 즉시 반환한다. 실제 완료는 MPI_Wait(또는 MPI_Waitall)에서 확인한다.

// 슬랩의 위/아래 halo를 양쪽 이웃과 동시에 교환
int up   = (rank + 1) % size;
int down = (rank - 1 + size) % size;

float *send_up, *send_down, *recv_up, *recv_down;
// ... 위/아래로 보낼 면 데이터 준비 (GPU 또는 host buffer) ...

MPI_Request req[4];
MPI_Irecv(recv_down, FACE, MPI_FLOAT, down, 1, MPI_COMM_WORLD, &req[0]);
MPI_Irecv(recv_up,   FACE, MPI_FLOAT, up,   2, MPI_COMM_WORLD, &req[1]);
MPI_Isend(send_down, FACE, MPI_FLOAT, down, 2, MPI_COMM_WORLD, &req[2]);
MPI_Isend(send_up,   FACE, MPI_FLOAT, up,   1, MPI_COMM_WORLD, &req[3]);

// ... 이 사이에 통신과 무관한 일을 할 수 있다 ...

MPI_Waitall(4, req, MPI_STATUSES_IGNORE);

이 패턴이 halo exchange의 표준 형태다. Irecv를 먼저 거는 것이 좋다 — 라이브러리에 "여기 받을 준비 됐어, 메시지 오면 바로 여기에 써" 하고 알려주면 추가 복사가 줄어든다. tag도 일관되게 쓴다(위로 보내는 건 항상 tag 1, 아래는 2 식). 위/아래 양쪽을 모두 동시에 비블로킹으로 걸면, 양방향 동시 진행으로 통신 시간이 약 절반.

20.5 통신과 계산 오버랩 — 진짜 가속의 핵심

여기까지가 정직한 halo exchange. 그런데 진짜 가속의 비결은 통신 시간을 계산 시간 뒤에 숨기는 것이다. MPI_Waitall이 끝날 때까지 GPU가 놀고 있으면 망한다. 다음 두 단계 분할이 답이다.

1. internal compute: halo가 필요 없는 슬랩 안쪽 점들을 먼저 갱신. 이건 통신과 독립이라 통신과 동시에 진행 가능.
2. boundary compute: halo가 필요한 슬랩 경계점은 MPI_Waitall 뒤에 갱신.

이걸 GPU에서 실현하려면 도구가 더 필요하다. (a) 핀드 메모리(pinned host memory) — 일반 malloc 메모리는 OS가 swap할 수 있어 DMA가 안전하지 않다. cudaMallocHost로 page-locked 메모리를 잡으면 DMA가 직접 가능해 PCIe 전송이 2~3배 빨라진다. (b) CUDA stream — 비동기 작업의 큐. 여러 stream에 큐를 걸면 작업들이 의존성에 따라 동시에 진행된다. (c) cudaMemcpyAsync — stream에 거는 비동기 복사.

// 두 stream: 통신용(halo)과 계산용(internal)
cudaStream_t s_comm, s_comp;
cudaStreamCreate(&s_comm);
cudaStreamCreate(&s_comp);

// host 핀드 staging buffer
float *h_send_up_pin, *h_recv_up_pin;
cudaMallocHost(&h_send_up_pin, FACE*sizeof(float));
cudaMallocHost(&h_recv_up_pin, FACE*sizeof(float));
// (down도 마찬가지)

void step(float* d_grid_in, float* d_grid_out) {
    // 1. 경계 면을 GPU -> 핀드로 비동기 복사 (s_comm)
    pack_face_kernel<<<..,s_comm>>>(d_grid_in, d_send_up_buf, /*top*/);
    pack_face_kernel<<<..,s_comm>>>(d_grid_in, d_send_down_buf, /*bot*/);
    cudaMemcpyAsync(h_send_up_pin,  d_send_up_buf,  FACE*4, cudaMemcpyDeviceToHost, s_comm);
    cudaMemcpyAsync(h_send_down_pin,d_send_down_buf,FACE*4, cudaMemcpyDeviceToHost, s_comm);

    // 2. 통신과 독립적인 슬랩 안쪽 갱신 시작 (s_comp)
    stencil_internal<<<..,s_comp>>>(d_grid_in, d_grid_out);

    // 3. host 복사 끝나면 MPI 비블로킹 송수신
    cudaStreamSynchronize(s_comm);
    MPI_Request req[4];
    MPI_Irecv(h_recv_down_pin, FACE, MPI_FLOAT, down, 1, MPI_COMM_WORLD, &req[0]);
    MPI_Irecv(h_recv_up_pin,   FACE, MPI_FLOAT, up,   2, MPI_COMM_WORLD, &req[1]);
    MPI_Isend(h_send_down_pin, FACE, MPI_FLOAT, down, 2, MPI_COMM_WORLD, &req[2]);
    MPI_Isend(h_send_up_pin,   FACE, MPI_FLOAT, up,   1, MPI_COMM_WORLD, &req[3]);

    // 4. 내부 계산은 그대로 진행 중. 둘이 겹친다.
    MPI_Waitall(4, req, MPI_STATUSES_IGNORE);

    // 5. 받은 halo를 GPU로 복사 후 경계 갱신
    cudaMemcpyAsync(d_recv_up_buf,  h_recv_up_pin,  FACE*4, cudaMemcpyHostToDevice, s_comm);
    cudaMemcpyAsync(d_recv_down_buf,h_recv_down_pin,FACE*4, cudaMemcpyHostToDevice, s_comm);
    cudaStreamSynchronize(s_comm);
    cudaStreamSynchronize(s_comp);   // 내부 갱신도 끝났는지 확인
    stencil_boundary<<<..>>>(d_grid_in, d_grid_out, d_recv_up_buf, d_recv_down_buf);
}

이 패턴의 정량 효과는 응용에 따라 다르지만 대체로 이렇다. 통신 시간 T_c, 내부 계산 시간 T_i, 경계 계산 시간 T_b일 때, 직렬은 T_c + T_i + T_b. 오버랩은 max(T_c, T_i) + T_b. T_c ≈ T_i면 약 1.5~1.8배 가속, T_i >> T_c면 거의 통신 비용이 사라진다.

20.6 collective 통신 — 한 줄로 끝나는 패턴들

점대점만 있으면 모든 분산 알고리즘을 짤 수 있지만, 자주 쓰는 패턴은 라이브러리가 미리 최적화해 둔 collective를 쓰는 게 빠르고 깔끔하다.

• MPI_Bcast: rank 0이 가진 데이터를 모든 rank에 뿌린다. 초기 파라미터 배포에 흔히 사용. 내부적으로는 트리 구조로 O(log P)에 끝난다.
• MPI_Scatter / MPI_Gather: rank 0이 큰 배열을 가지고 있고 균등하게 나눠 뿌리는(혹은 거꾸로 모으는) 패턴. 입력 분배, 결과 수집에 쓴다.
• MPI_Reduce / MPI_Allreduce: 모든 rank의 partial sum을 합산. 학습에서 gradient 평균에 매번 쓰인다. Allreduce는 결과를 모든 rank에 다시 뿌린다. 트리 + 링 결합으로 O(2(P-1)/P · M)의 대역폭이 이론치.
• MPI_Barrier: 모든 rank가 도착할 때까지 대기. 사실 잘 디자인된 코드에서는 거의 안 써도 된다. 디버깅용으로는 가끔 쓴다.

딥러닝 데이터 병렬 학습에서는 매 step마다 모든 GPU의 gradient를 Allreduce로 평균낸다. 이 한 줄이 분산 학습의 심장이다. 점대점으로 같은 일을 직접 짜면 P 노드에서 P-1 단계의 통신이 필요하지만, ring-allreduce 알고리즘은 P-1 단계로 끝난다(메시지 크기가 작아짐). collective의 장점은 라이브러리가 이런 알고리즘을 이미 골라준다는 점.

20.7 CUDA-aware MPI — host staging을 건너뛰자

20.5에서 본 GPU↔host↔MPI 왕복은 본질적으로 군더더기다. 데이터는 어차피 GPU에 있고 다음에도 GPU에서 쓸 건데, 왜 host를 거쳐야 하는가? CUDA-aware MPI(OpenMPI/MVAPICH 등이 지원)는 GPU device pointer를 MPI 함수에 직접 넘길 수 있게 해준다.

// CUDA-aware MPI: GPU 포인터를 그냥 넘긴다
float* d_send_up;   // device pointer
float* d_recv_up;   // device pointer

MPI_Isend(d_send_up, FACE, MPI_FLOAT, up, 1, MPI_COMM_WORLD, &req[0]);
MPI_Irecv(d_recv_up, FACE, MPI_FLOAT, up, 2, MPI_COMM_WORLD, &req[1]);

구현 측에서 두 가지 일이 일어난다. (a) 같은 노드 내라면 NVLink/PCIe peer-to-peer로 host를 거치지 않고 GPU↔GPU 직접 복사. (b) 노드 간이라면 GPUDirect RDMA로 NIC가 GPU 메모리를 직접 DMA해 host RAM을 거치지 않는다. 이 둘이 결합되면 latency가 수십 마이크로초에서 수 마이크로초로 떨어지고, 대역폭은 PCIe 한계에 가깝게 올라간다.

정량적으로, 64KB halo 면 한 번 교환에 host staging 방식이 ~50µs 수준이라면 GPUDirect RDMA는 ~10µs 수준. 1000 step짜리 시뮬레이션이면 통신만 따져 40초 절약. 강력한 스텐실에서는 step time 자체가 ms 단위라, 통신 hiding과 GPUDirect를 둘 다 적용하면 노드 4대에서 strong scaling이 거의 P에 가깝게 나오는 경우가 흔하다.

20.8 정리

한 GPU에서 노드 클러스터로 옮기면 통신이 dominant 비용이 된다. 우리는 이 장에서 다음 도구들을 차례로 꺼냈다 — 도메인 분할로 일을 나눈다, halo exchange로 경계를 동기화한다, 비블로킹 + 두 단계 분할로 통신을 계산 뒤에 숨긴다, CUDA-aware MPI와 GPUDirect로 host staging까지 없앤다. 한 단계씩 통신 비용을 깎아 나가는 사고 흐름은 17~18장의 단일-GPU 최적화와 본질적으로 같다. 단지 무대가 클러스터로 옮겨졌을 뿐.

여기서 강조해 둘 한 가지. 멀티노드 분산은 디버깅이 잔혹하다. 한 rank만 deadlock에 빠져도 클러스터 전체가 정지하고, 어디서 멈췄는지 알아내기가 어렵다. MPI_Barrier로 단계마다 동기화하고, rank마다 로그 prefix를 다르게 찍고, 작은 P=2 케이스부터 단계적으로 확장하는 습관이 시간을 아낀다. 그리고 가장 중요한 첫 디버깅 도구 — 코드의 통신을 도식으로 그려보는 것이다. 누가 누구에게 무엇을 보내는지를 그림으로 그릴 수 없으면 코드로 짤 수도 없다.