[Multi-Core & GPU Programming] CUDA - Multi GPU

"CUDA - Multi GPU"

[Multi GPU]

▣ Multi GPU 사용 환경

> Work Station / Desktop

· 일반 PC에서 여러 GPU 사용(1 CPU + 1 Main Memory + 여러 GPU)

> Server

· Desktop보다 훨씬 많은 GPU  장착 가능

> 여러 Node로 구성된 Cluster

· GPU가 여러 개 들어간 Node(=Server)가 여러 대 있고, 이 Node(=Server)들이 Network로 연결

[Single-Node Multi GPU]

· 하나의 Computer 또는 하나의 Server 안에 GPU가 여러 개 있는 구조

· cudaSetDevice( ) : 앞으로 어떤 GPU를 사용할 것인지 설정 / cudaMemcpy : CPU-GPU, GPU-GPU 간 데이터를 복사할 때 사용

▣ cudaSetDevice( )

· cudaSetDevice(0) : 현재 CUDA 작업의 대상 = GPU 0 & Kernel 실행 시 GPU 0 에서 실행

    → 현재 CPU Thread가 CUDA 명령을 보낼 대상 GPU를 바꾸는 함수

· cudaMemcpyAsync(…) : 복사를 요청한 후 CPU가 해당 복사가 끝날 때까지 기다리지 않고 다음 명령어 실행

    → GPU 0은 아직 데이터를 복사하고 있고, 동시에 GPU 1은 다른 Kernel을 실행 가능

▣ Unified Addressing

· CPU와 여러 GPU가 각각 자신만의 Virtual Memory Address Region 보유 가능

→ CPU Memory 주소, GPU 0 Memory 주소, GPU 1 Memory 주소를 하나의 통합된 주소 체계 안에서 구분 가능

** 모든 GPU가 자동으로 서로의 Memory를 자유롭게 Read/Write X

▣ GPU-GPU Communication

· cudaDeviceEnablePeerAccess(peer_device, 0) : 현재 GPU가 다른 "peer_device" GPU의 메모리 주소에 접근할 수 있도록 허용

· cudaDeviceCanAccessPeer(&accessible, dev_X, dev_Y) : "dev_X"가 "dev_Y"의 Memory에 접근 가능한지 확인

∴ GPU끼리 직접 데이터를 주고 받으려면 먼저 GPU들이 서로 접근 가능한지 확인 후, 가능하다면 Peer Access를 활성화

· cudaMemcpyPeerAsync(…) : GPU끼리 비동기적으로 데이터를 복사하는 함수

    → "src_dev" GPU의 "src_addr"에서 데이터를 읽어서, "dst_dev" GPU의 "dst_addr"로 "num_bytes"만큼 복사

    → 마지막 인자인 "stream"은 복사 작업을 어느 CUDA Stream에 넣을지 결정

· 현재 성능은 데이터가 출발하는 Source GPU에 속한 Stream을 사용할 때 가장 극대화되며, 비동기 방식이 아닌 동기 버전의 함수도 존재

· Peer-Access 활성화 : 데이터가 CPU Memory를 거치지 않고 가장 짧은 PCIe 경로를 통해 직접 전송

· Peer-Access 비활성화 : CUDA Driver가 데이터를 CPU Memory로 먼저 옮긴 후 목적지 GPU로 전달 → 속도↓

▣ Consider Topology

· 여러 GPU가 모두 같은 방식으로 연결 X(같은 PCIe Switch / 다른 PCIe Switch)

> 정상 동작

> Conflict 발생

· 동시에 같은 PCIe 경로 사용 → Conflict → 속도↓

▣ Stream & Event

· CUDA Stream과 Event는 Device별로 존재(GPU 0에서 만든 Stream = GPU 0에 귀속 / GPU 1에서 만든 Stream = GPU 1에 귀속)

· 각 GPU는 자기만의 Default Stream을 보유(GPU 0의 Default Stream ≠ GPU 1의 Default Stream)

· Kernel : 해당 Stream이 속한 GPU가 현재 활성화된 상태일 때만 실행 가능

· Memcpy : 현재 활성화된 GPU와 상관없이 어느 Stream에든 명령 가능

· Event : Event를 기록하는 것은 Event, Stream이 모두 현재 활성화된 GPU에 속한 상태일 때만 가능

· 동기화/Query : 다른 Device에 속한 Stream/Event를 대상으로 수행해도 무방

> Event

· Kernel 실행 시간을 측정할 때 사용(어떤 작업이 끝날 때까지 기다리는 용도로 사용)

→ 같은 Stream내의 다른 Kernel에서 실행되는 데이터에 의존성이 있어 동기화가 필요한 경우 사용

    (∵ __synctrheads( )는 같은 Kernel, 같은 Block 내부 Thread 동기화)

→ 같은 GPU 내의 다른 Stream에서 실행되는 데이터에 의존성이 있어 동기화가 필요한 경우 사용

→ 다른 GPU에서 실행되는 데이터에 의존성이 있어 동기화가 필요한 경우 사용

① "stream1"에 Kernel을 실행시키고, 그 뒤에 "event1"을 기록

② "event1"은 "stream1"에서 앞의 작업들이 여기까지 끝났다"는 표시

③ "cudaEventSynchronize(event1)" 호출 시 CPU는 "event1"이 완료될 때까지 대기

→ Event : Multi-GPU 또는 Multi-Stream 환경에서 작업 완료 시점을 표시하고 기다리는 도구

> Example 1

① streamA & eventA : GPU 0에 귀속 / streamB & eventB : GPU 1에 귀속

② Kernel : streamB 실행(현재 Device = GPU 1)

③ Event : eventB 기록(현재 Device = GPU 1)

④ 동기화 : 다른 Device여도 상관 X

→ 정상

> Example 2

① streamA & eventA : GPU 0에 귀속 / streamB & eventB : GPU 1에 귀속

② Kernel : streamA 실행(현재 Device = GPU 1)

③ Event : eventB 기록(현재 Device = GPU 1)

④ 동기화 : 다른 Device여도 상관 X

→ 비정상(∵ Kernel Launch 규칙 위배)

∴ Kernel을 실행하기 전에 "cudaSetDevice(0);" 추가 필요

> Example 3

① streamA & eventA : GPU 0에 귀속 / streamB & eventB : GPU 1에 귀속

② Kernel : streamB 실행(현재 Device = GPU 1)

③ Event : eventA 기록(현재 Device = GPU 1)

→ 비정상(∵ Event Record 규칙 위배)

∴ "cudaEventRecord(eventB, streamB);"로 변경 필요

> Example 4

① streamA & eventA : GPU 0에 귀속 / streamB & eventB : GPU 1에 귀속

② Kernel : streamB 실행(현재 Device = GPU 1)

③ Event : eventB 기록(현재 Device = GPU 1)

④ Device 변경(GPU 1 → GPU 0)

⑤ 동기화 : 다른 Device여도 상관 X

⑥ Kernel : streamA 실행(현재 Device = GPU 0)

→ 정상

> Example 5

① GPU 0 선택 후, GPU 0 Memory에 "d_A" 할당

② GPU 1이 GPU 0의 Memory에 접근할 수 있는지 확인

③ 접근 가능한 경우 현재 Device를 GPU 1로 변경 후, GPU 0 Memory "d_A"에 접근(Kernel은 GPU 1에서 실행)

[Multi-Node Multi GPU]

· 병렬 Programming에서 Programming Model은 여러 계산 주체가 데이터를 어떻게 주고받는지 정하는 방식

  ① Shared Memory Model : 여러 CPU 또는 Core가 같은 Memory 공간 공유

  ② Message-Passing Model : 각 Process가 자기 Memory를 따로 가지고 있고, 데이터를 주고 받으려면 명시적으로 메세지 전송

▣ Distributed Memory Model

· CPU마다 자기 Memory가 따로 존재 & 두 CPU는 Network로 연결

→ 데이터를 주고 받으려면 Message 전달 후 Network를 통해 복사

> Shared Memory Model

· 여러 CPU가 같은 Memory에 있는 "global_sum"에 접근 가능

△ 사용 편리

△ 같은 Memory를 공유하므로 Programming이 직관적

▼ 보통 많은 Node(100 Core 이상)로 확장 불가

> Distributed Memory Model

· 각 CPU가 자기 Memory를 따로 가지고 있으므로, 한 쪽 Memory에 있는 "global_sum" 접근 불가

→ 데이터를 보내려면 Message 통신을 통해서 복사

△ 여러 Server를 Network로 연결해서 더 큰 시스템으로 확장 가능

▼ 복사 과정이 필요하여 Shared Memory Model보다 속도↓

▣ Message Passing

· 데이터를 공유하기 위해 명시적으로 Message를 주고받는 방식

· Core들은 Message를 서로 주고받는 Handshaking Message를 통해 데이터를 공유

· 데이터를 보내는 경우 : Send( ) / 데이터를 받는 경우 : Recv( ) → 두 함수 호츨은 반드시 짝이 맞아야 함

▣ MPI

· MPI_Init : MPI 환경 시작

· MPI_Comm_size : 전체 Process 개수 계산

· MPI_COMM_WORLD : 모든 MPI Process를 포함하는 기본 Communicator

· MPI_Comm_rank : 현재 Process의 번호 계산

· MPI_Get_processor_name : 현재 Process가 실행 중인 Host 이름 계산

· MPI_Finalize( ) : MPI 종료

· Compile : mpig++ main.cc / 실행 : mpirun -n 4 ./a.out (Process를 4개 실행)

· 총 4개의 MPI Process가 실행되었고, 각각 자신의 Process 번호를 출력

> Send( )

· smessage : 전송할 데이터가 담긴 Buffer

· count : 전송할 데이터 요소의 개수
· datatype : 전송할 데이터의 Type(Ex. MPI_INT, MPI_FLOAT 등)

· dset : 데이터를 받을 목적지 Process의 Rank

· tag : 수신자가 Message를 구분할 수 있게 해주는 Tag

· comm : 통신에 사용되는 Communicator(기본값 = MPI_COMM_WORLD)

> Recv( )

· rmessage : 수신할 데이터를 저장할 Buffer
· count : 수신할 데이터 요소의 개수

· datatype : 수신할 데이터의 Type

· source : 데이터를 보낸 Process의 Rank

· tag : Message를 구분하기 위한 Tag

· comm : 통신 Communicator

· status : 수신된 Message의 상태 정보를 담는 자료 구죠

▣ MPI Works For GPU

· GPU에서 MPI를 이용해 서로 다른 Node의 GPU 데이터를 사용 가능

· 예전에는 GPU Memory에 있는 데이터를 MPI로 바로 보내기 불가

    (GPU Memory → CPU Host Memory → MPI Send → MPI Recv → CPU Host Memory → GPU Memory)

▣ MPI Data Type

[Reference]

· 14_multi_gpu_mgp_2026 (MGP) - Yongjun Park