SHPC 17.5 - Multiple GPUs (CUDA)

일반적인 GPU 시스템의 구성

• 12 CPU + 18 GPU
• PCIe를 통해 CPU와 GPU를 연결
• GPU와 CPU를 보조 장치로 사용

CUDA 프로그램 구성

• 호스트 프로그램
  • 평범한 C/C++ 프로그래밍 언어로 작성
  • 메인 프로세서 (CPU) 에서 실행
  • CUDA API를 사용해 디바이스에 작업을 요청
    • 계산, 데이터 전송, 동기화 등
• 커널
  • CUDA C/C++ 확장 언어로 작성
  • 디바이스에서 실행되는 코드의 기본 단위
• 호스트 프로그램과 커널은 동시에 실행 가능
  • CPU와 GPU가 동시에 작업을 수행할 수 있음
  • 여러 GPU를 동시에 사용하는 것도 가능

일반적인 CUDA 프로그램 실행 과정

CUDA Multi-GPU

• 다수의 GPU를 다루기 위한 방법
  • 한 CPU 스레드가 여러 GPU를 관리하는 방법 (1thread - NGPU)
  • 한 CPU 스레드가 한 GPU를 관리하는 방법 (1thread - 1GPU)
    • Pthread, OpenMP 등으로 여러 CPU 스레드를 생성
    • MPI로 여러 프로세스 생성
  • 상황에 따라 적절한 방식 선택
• 프로그래밍 방법

1. 사용 가능한 GPU 개수 확인
2. CUDA API 의 대상이 되는 GPU 지정 cudaSetDevice(int device)
   • cudaSetDevice(i) 이후 호출되는 모든 작업은 i번째 GPU에서 실행됨
   • 스트림 및 이벤트도 마찬가지
     • GPU마다 각각 스트림과 이벤트를 생성해 주어야 함

• 예시
• Async를 하는 이유는 CPU와 동시 실행을 하기 위함
• 커널은 stream을 꼭 넣어줘야 한다
  • Stream: GPU에 작업을 넣는 큐
  • GPU마다 하나의 stream을 만들어 줘야 한다

다수의 GPU로 행렬곱 하기

• 결과 C 행렬을 각 GPU가 나눠가진다고 생각해 보자
• 먼저 타임라인을 그리는 것이 중요!!

다수의 GPU로 작업하는 방법

• 하나의 CPU 스레드가 다수의 GPU를 맡아서 관리

• 모든 CUDA API를 비동기로 호출해야함

  • 동기 API로 호출한다면?

• 각 GPU에 작업을 순서에 맞게 비동기로 실행시켜 놓은 뒤 한번에 동기화

• 일반적인 형태

다수의 GPU로 행렬곱 하기

• 각 GPU가 나누어 가질 행렬 인덱스 계산
  • i번 GPU는 C의 [Mbegin[i], Mend[i]] 행을 담당해 계산

1. GPU마다 스트림 및 이벤트 생성
   • 이벤트는 필요 없음
   • 더블/트리플 버퍼링까지 적용하려면 이벤트 필요 < 설명 X
2. GPU마다 작업에 필요한 메모리 할당
   • 필요한 만큼만 할당하기
3. GPU마다 입력 데이터를 비동기로 전송

• 호스트 메모리가 Pinned 메모리가 아닌 경우 동기로 작동함에 주의!!!!

4. GPU마다 커널 실행
5. GPU마다 출력 데이터를 비동기로 전송
6. 모든 GPU작업이 끝날 때까지 대기
7. 생성/할당한 자원 해제

비동기 메모리 복사 유의점

• 메모리 복사의 경우 비동기 API를 사용해도 실제로 비동기로 동작하지 않는 경우가 있음
  • cudaMemcpyAsync(…)
• 비동기 API가 동기로 동작하는 두 가지 경우
  1. 디바이스 메모리 ←> pageable 호스트 메모리 (Pinned가 아닌 경우)
  2. 호스트 메모리 ←> 호스트 메모리
• 비동기 API로 데이터 전송할 때는 호스트 메모리가 pinned 메모리가 맞는지를 반드시 체크해야 함
  • cudaHostAlloc(), cudaMallocHost() 로 할당 가능
  • cudaFreeHost() 로 해제

다수의 CPU 스레드 활용하기

• 자원 할당/해제는 하나의 스레드가 맡아서 해도 무방함
• GPU 하나당 하나의 CPU 스레드를 생성
  • 스레드별로 하나의 GPU의 데이터 전송 및 커널 실행을 맡아 실행
  • 각 스레드 내 CUDA API 호출은 비동기일 필요가 없음
    • 동기 메모리 복사 등 사용 가능
• 마지막에 sync와 join을 꼭 해줘야 한다
• 아래는 예시

cudaStreamCreate도 setdevice를 하고 호출해야 했다