NCCL 설치 및 행렬곱 예시 (Scatter, Broadcast, Gatter)

이전글

2022.07.22 - [Parallel Computing/개념] - NCCL 개념 및 Ring 기반 집합 통신 최적화

이전 글에서 NCCL[1]에 대한 기본 개념에 대하여 정리하였다. 이번 글에서는 NCCL을 직접 설치해보고 NCCL API를 사용하여 행렬곱을 구현해 보고자 한다.

 

 

 

NCCL 설치

NCCL[Nickel]은 NVIDIA에서 배포한 GPU간 최적화된 집합 통신(collective communicatio) library이다. 여러 GPU간의 데이터 전송 관련 primitives를 구현한 것으로, 가속화(최적화)된 GPU간의 데이터 통신 구현에만 초점을 맞춘 library이다. 이를 통해 NVIDIA는 쉬운 multi-GPU 애플리케이션 개발을 지원하고자 한다. 

 

NCCL은 library이기에 OS와 driver에 맞는 라이브러리 파일(.so, 동적 라이브러리)과 헤더 파일만 다운로드 받으면 된다. NCCL 라이브러리 파일과 헤더 파일은 다음에서 다운로드 받을 수 있다.

 

NCCL 테스트 환경은 CentOS 8.5, CUDA 11.7이다. 설치를 위해 CUDA 11.7 버전용 NCCL [O/S agnostic local installer]를 다운로드 받으면, ".txz" 압축파일이 다운로드 된다. 이 압축파일 내부에는 NCCL 사용을 위한 라이브러리 파일들과 헤더 파일들이 있다. 이들을 원하는 위치에 압축해제하고 환경변수 설정만 완료하면 된다. 나는 환경변수를 설정할 필요 없게 CUDA 설치 환경에 NCCL 라이브러리와 헤더를 설치하였다.

 

 

 

2개의 GPU를 이용한 행렬곱 구현

행렬 A, B에 대하여 행렬 C = AB를 구하고자 한다. 다음 Fig 1.과 같이 2개의 GPU를 이용한 행렬곱을 구현할 수 있다.

 

FIg 1. GPU 2개를 이용한 행렬곱 구현

 

Fig 1.에서 볼 수 있듯이, 2개의 GPU 각각에서 C 행렬 반씩 계산하고, 그 값을 합치고자 한다. 이를 위해 행렬 A를 GPU1과 GPU2에 반씩 scatter한다. 행렬 B의 경우에는 행렬 B가 모두 필요하기에 B를 GPU1, GPU2에 broadcast한다. GPU1는 행렬 C의 윗쪽 반을, GPU2는 행렬 C의 아랫쪽 반을 계산한다. 최종 계산된 결과를 gather하여 행렬 C를 완성한다. 

 

다음과 같은 구현 시, 개별 GPU당 메모리 사용량 및 연산 시간이 줄어들 수 있다. 예를 들어 M=K=N일 시, 개별 GPU 당 메모리 사용량은 2/3가 되며 계산 시간 또한 1/2로 줄어든다. 데이터 통신량은 4/3로 늘어나는 것을 알 수 있다. (B를 두번 전송하여야 한다) 

 

 

 

구현

다음 github[3]에 GPU2개를 이용한 행렬곱 구현을 올려두었다. 

GitHub - jhson989/nccl-matmul-tutorial: NCCL example for a matrix multiplication application in a single node

 

시스템 메모리에서 GPU 2개에 데이터(행렬 A, B)를 전송하고, GPU2개에서 다시 시스템 메모리로 데이터(행렬 C)를 전송할 때 NCCL을 사용하여 구현하였다. NCCL은 GPU간 데이터 통신을 위한 라이브러리라고 하였지만, CPU-GPU간 데이터 통신(점대점 통신, 집합 통신)도 지원한다. 다만 CPU-GPU간 데이터 통신은 cudaMemcpy를 사용해서 직접 구현해도 되지만(사실 성능 상의 이점도 없을 것 같다), NCCL 공부 겸 NCCL을 이용하여 구현해 보았다.

 

행렬 A Scatter

NCCL은 Scatter 함수를 지원하지 않고, 점대점 통신을 통해서 구현하도록 한다[4].

 

// Point-to-point communication을 이용한 Scatter 구현
ncclErrChk( ncclGroupStart() );
for (int i=0; i<num_devices; i++) {
    ncclErrChk( ncclSend(A.data()+N*N/num_devices*(i), N*N/num_devices, ncclFloat32, i, comms[i], streams[i]) );
    ncclErrChk( ncclRecv(d_A[i], N*N/num_devices, ncclFloat32, i, comms[i], streams[i]) );
}
ncclErrChk( ncclGroupEnd() );

 

ncclSend 함수를 통해 시스템 메모리에 저장된 A 행렬을 2개의 GPU 메모리 영역(d_A[0], d_A[1])에 반반씩 전송한다. ncclReav 함수를 통해 GPU에게는 전송받도록 명령한다. 

 

이때 Group call[7](GroupStart, GroupEnd 함수)을 통해 ncclSend()들과 ncclRecv()들이 묶여있는 것을 확인할 수 있다. Group call은 여러 개의 통신이 동시에 호출될 경우 하나로 합쳐주는 통신 명령어로 deadlock을 방지하는 효과가 있다. ncclSend와 ncclRecv는 완료될 때까지 GPU들이 blocking되는데, 만약 GPU0에서 GPU1에 데이터를 보내는 동시에 GPU1에서도 GPU0으로 데이터를 보내면 deadlock이 발생할 수 있다. Group call은 이러한 deadlock을 방지하는 목적으로 사용되며, NCCL 통신을 쓸 때는 항상 사용해야 하는 것으로 보인다. (사실 Group call을 쓰지 않고 실행해보려고 하였지만 error가 발생하며 실행이 되지 않았다. 그래서 안쓰고 사용해도 되는 지는 확인하지 못했다.)

 

행렬 B Broadcast

NCCL은 집합 통신을 위해 broadcast 함수를 지원한다[5].

 

// NCCL을 통한 Broadcast
ncclErrChk( ncclGroupStart() );
for (int i=0; i<num_devices; i++) {
    ncclErrChk( ncclBroadcast(B.data(), d_B[i], N*N, ncclFloat32, 0, comms[i], streams[i]) );
}
ncclErrChk( ncclGroupEnd() );
sync_all();

 

ncclBroadcast 함수를 통해 시스템 메모리에 저장된 B 행렬을 2개의 GPU 메모리 영역(d_B[0], d_B[1])에 모두 전송한다. 

 

행렬 C Gather

NCCL은 Gather 함수를 지원하지 않고, 점대점 통신을 통해서 구현하도록 한다[6].

 

// Point-to-point communication을 통한 Gather 구현
ncclErrChk( ncclGroupStart() );
for (int i=0; i<num_devices; i++) {
    ncclErrChk( ncclSend(d_C[i], N*N/num_devices, ncclFloat32, i, comms[i], streams[i]) );
    ncclErrChk( ncclRecv(C.data()+N*N/num_devices*(i), N*N/num_devices, ncclFloat32, i, comms[i], streams[i]) );
}
ncclErrChk( ncclGroupEnd() );

 

ncclSend 함수를 통해 2개의 GPU 메모리 영역(d_C[0], d_C[1])에 나뉘어 저장된 행렬 C를 시스템 메모리 C에 저장한다. ncclReav 함수를 통해 CPU에게는 전송받도록 명령한다.

 

 

 

Reference

[1] https://docs.nvidia.com/deeplearning/nccl/index.html

[2] https://developer.nvidia.com/nccl/nccl-download

[3] https://github.com/jhson989/nccl-matmul-tutorial

[4] https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/p2p.html#one-to-all-scatter

[5] https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/collectives.html#broadcast

[6] https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/p2p.html#all-to-one-gather

[7] https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/groups.html